C++ GPGU Tutorial

Danke für den Hinweis, mir war überhaubt nicht bewusst dass unter CUDA normalerweise keine fehlermeldungen ausgegeben werden.
Ich habe mitlerweile eine kleine sicherheitsabfrage eingebaut und dass Array hat jetzt eine größe von 10k.
Die richtige Antwort ist übrigends 983, denn 1/983 hat eine 1866 stellige Periode :freak:
 
Du solltest mit den automatischen Arrays etwas aufpassen, da diese im CUDA lokalen Speicher abgelegt werden. Dieser ist atm bei Compute Capability 2.0 bis 5.0 auf 512 KB pro Thread begrenzt. Zudem kostet dir das Speicherplatz ohne Ende: Ein lokales int Array[10 000] kostet dir bei der maximalen Occupancy von 64 insgesamt auf einer Titan 1.2 GB an Speicher, wodurch das ganze auch dementsprechend Cache unfreundlich ist.
 
Das die Arrays im lok speicher abgelegt werden, sofern nicht anders definiert ist mir bewusst, jedoch war ich der Ansicht, dass sie klein genug dafür sind. Andernfalls hätte ich sie halt im shared bzw Global memory erstellt.
Bei Cpu`s habe ich z.T. vektoren / stacks gearbeitet jedoch habe ich noch nichts darüber gelesen wie weit das auch auf GPU`s funktioniert, falls du da genauere Infos hast währe ich dir dankbar. In dem Udacity kurs den ich langsam aber sicher durchgehe kam das noch nicht dran Das Nvidia tutorial habe ich schon länger nicht benutzt da es für mich z.Z zu viele informationen enthält und ich mit dem Udacity kurs deutlich entspannter lerne.
 
Bekommst du denn irgendwelche Fehlermeldungen nachdem du das Errorchecking eingebaut hast, bzw. hast du es auch richtig verwendet - cudaDeviceSynchronize() zwischen Kernelaufruf und cudaGetLastError()?
Da ich aktuell gerade unterwegs bin, könnte ich erst Ende der Woche deinen Code nehmen und ausführen.
 
Zuletzt bearbeitet:
Da mein altes Programm läuft mal eine ganz andere Frage, für eine neuronales Netzwerk brauche ich viele vier gewinnt Partien.
Diese können vollkommen wahllos gespielt werden, weshalb ich das ganze über CUDA lösen würde.
Die Partien würde ich von jedem thread als Array aus 43 shorts (42 züge + ergebnis) zurückgeben lassen.
Eine GTX 470 besitzt ja 448 SM`s und pro SM X cuda cores.
Bei 1024 threads/core * x cores/SM * 448 SM`s * 2 byte/short * 43 shorts = 39.452.672 * x threads
sehr viele threads, wie viele cores es pro SM gibt konnte ich leider nicht rausfinden. Könnte ich da speicherprobleme bekommen?
bei grob 100 bytes / thread macht das bei 1 GB speicher insgesamt ~10.000.000 threads die ich gleichzeitig nutzen könnte also onehin mehr als ich maximal erstellen kann oder habe ich mich irgendwo verrechnet?
 
Das die Arrays im lok speicher abgelegt werden, sofern nicht anders definiert ist mir bewusst, jedoch war ich der Ansicht, dass sie klein genug dafür sind. Andernfalls hätte ich sie halt im shared bzw Global memory erstellt.

Das ändert doch das Problem nicht. Im Shared Memory ruinieren sie dir komplett deine Occupancy und als einfaches globales Array belegen sie genauso viel Speicherplatz. Das Ziel sollte es eher sein sich zu überlegen, wie man solche Arrays bei dem Algorithmus komplett vermeiden kann - ist es aus algorithmischen Gründen nur schlecht möglich, so handelt es sich wahrscheinlich um einen Algorithmus der sich eher weniger für GPUs eignet.



Bei Cpu`s habe ich z.T. vektoren / stacks gearbeitet jedoch habe ich noch nichts darüber gelesen wie weit das auch auf GPU`s funktioniert, falls du da genauere Infos hast währe ich dir dankbar.
Afaik gibt es atm keine Lib, welche solche dynamischen Datenstrukturen innerhalb eines Kernels unterstützt (habe mich damit aber auch noch nicht beschäftigt, eben weil es relativ unüblich ist.). Du kannst dir aber selbst eine solche dynamische Datenstruktur schreiben:
-Entweder mit malloc im Kernel, dafür musst du aber die Heap-Size deiner GPU entsprechend hoch einstellen
-Oder du schreibst dir selbst eine Routine welche in einem globalen Array freie Blöcke sucht, alloziert und anschließend wieder freigibt.

Da Malloc prinzipiell hardwarebeschleunigt ist, nehme ich an, dass die erste Möglichkeit schneller ist.


Eine GTX 470 besitzt ja 448 SM`s und pro SM X cuda cores.
Bei 1024 threads/core * x cores/SM * 448 SM`s * 2 byte/short * 43 shorts = 39.452.672 * x threads

Das ganze berechnet sich zu:
Threads pro SM * Anzahl an SMs
Bei maximaler Occupancy erreichst du:
1536 *14 Threads = 21 504 Threads

Also kann deine GPU zu einem Zeitpunkt eine Zahl von 21504 Threads, welche auch eine Wave genannt wird, gleichzeitig bearbeiten. Um eine etwas unterschiedliche Laufzeit zwischen den Threadblocks zu kompensieren und um den Overhead des Datentransfers und des Kernelstarts bei so einfachen kurz andauernden Kernels zu reduzieren empfiehlt es sich dass man bei einem Kernelstart eine Threadzahl startet welche am besten ein größeres ganzzahliges Vielfaches von einer Wave ist. Für ein paar schöne bunte Bilder dazu siehe folgende Slides: http://on-demand.gputechconf.com/gtc/2012/presentations/S0514-GTC2012-GPU-Performance-Analysis.pdf

Bei den geplanten 86 Bytes pro Thread hast du in etwa 2 Mbyte pro Wave, wodurch du maximal mehrere 100 Waves starten kannst, was ausreichend sein sollte.

Ich sehe bei dem Algorithmus eher ein anderes performancetechnisches Problem: Das Spielen einer 4-gewinnt Partie ist relativ einfach und benötigt kaum Rechenleistung, wodurch das Kernel selbst wahrscheinlich speicherbandbreiten limitiert ist. Da du die Daten allerdings noch auf dem Host haben willst musst du dir sie über den PCI-E ziehen, wodurch das ganze insgesamt eben sehr stark durch dessen Bandbreite limitiert wird. Deshalb ist der Algorithmus wahrscheinlich ebenfalls eher schlecht für GPGPU.
 
Danke, da habe ich mich ja ordentlich verrechnet.

Ich werde die Berechnungen dann vollständig auf der CPU durchführen, da habe ich dann auch genug RAM damit selbst meine ineffizienteren Anfangsalgorithmen vernünftig drauf laufen.
Bzgl. des alten Problems werde ich mich mal genauer über malloc informieren und gucken ob ich da was zurecht basteln kann.
 
Zurück
Oben