Ας πάμε πίσω στο 2003, όταν η Intel και η AMD βρίσκονταν σε κοινή κούρσα για την εύρεση του πιο ισχυρού επεξεργαστή. Μέσα σε λίγα μόλις χρόνια, ως αποτέλεσμα αυτού του αγώνα, οι ταχύτητες του ρολογιού έχουν αυξηθεί σημαντικά, ειδικά μετά την κυκλοφορία του Intel Pentium 4.

Όμως ο αγώνας πλησίαζε γρήγορα το όριο. Μετά από ένα κύμα τεράστιων αυξήσεων στις ταχύτητες ρολογιού (μεταξύ 2001 και 2003, η ταχύτητα ρολογιού του Pentium 4 διπλασιάστηκε από 1,5 σε 3 GHz), οι χρήστες έπρεπε να είναι ικανοποιημένοι με τα δέκατα των gigahertz που οι κατασκευαστές μπόρεσαν να αποσπάσουν (από το 2003 έως το 2005 , οι ταχύτητες ρολογιού αυξήθηκαν από μόνο 3 σε 3,8 GHz).

Ακόμη και αρχιτεκτονικές βελτιστοποιημένες για υψηλές συχνότητες ρολογιού, όπως το Prescott, άρχισαν να αντιμετωπίζουν δυσκολίες, και αυτή τη φορά όχι μόνο αυτές παραγωγής. Οι κατασκευαστές τσιπ απλώς έτρεξαν στους νόμους της φυσικής. Ορισμένοι αναλυτές μάλιστα προέβλεψαν ότι ο νόμος του Μουρ θα έπαυε να ισχύει. Αυτό όμως δεν συνέβη. Η αρχική έννοια του νόμου συχνά παραμορφώνεται, αλλά αφορά τον αριθμό των τρανζίστορ στην επιφάνεια του πυρήνα του πυριτίου. Για μεγάλο χρονικό διάστημα, η αύξηση του αριθμού των τρανζίστορ σε μια CPU συνοδεύτηκε από αντίστοιχη αύξηση της απόδοσης - η οποία οδήγησε σε παραμόρφωση του νοήματος. Στη συνέχεια όμως η κατάσταση έγινε πιο περίπλοκη. Οι προγραμματιστές της αρχιτεκτονικής της CPU προσέγγισαν το νόμο της μείωσης της ανάπτυξης: ο αριθμός των τρανζίστορ που έπρεπε να προστεθούν για την απαιτούμενη αύξηση της απόδοσης γινόταν ολοένα και μεγαλύτερος, οδηγώντας σε αδιέξοδο.

Ενώ οι κατασκευαστές CPU έσκιζαν τα μαλλιά τους προσπαθώντας να βρουν μια λύση στα προβλήματά τους, οι κατασκευαστές GPU συνέχισαν να επωφελούνται εντυπωσιακά από τα οφέλη του νόμου του Moore.

Γιατί δεν έφτασαν στο ίδιο αδιέξοδο με τους προγραμματιστές της αρχιτεκτονικής CPU; Ο λόγος είναι πολύ απλός: οι κεντρικοί επεξεργαστές έχουν σχεδιαστεί για να επιτυγχάνουν τη μέγιστη απόδοση σε μια ροή εντολών που επεξεργάζονται διάφορα δεδομένα (ακέραιους και αριθμούς κινητής υποδιαστολής), εκτελούν τυχαία πρόσβαση στη μνήμη κ.λπ. Μέχρι τώρα, οι προγραμματιστές προσπαθούν να παρέχουν μεγαλύτερο παραλληλισμό εντολών - δηλαδή να εκτελούν όσο το δυνατόν περισσότερες εντολές παράλληλα. Για παράδειγμα, με το Pentium, εμφανίστηκε η υπερκλιμακωτή εκτέλεση, όταν υπό ορισμένες συνθήκες ήταν δυνατή η εκτέλεση δύο εντολών ανά κύκλο ρολογιού. Το Pentium Pro έλαβε ακατάλληλη εκτέλεση εντολών, γεγονός που κατέστησε δυνατή τη βελτιστοποίηση της λειτουργίας των υπολογιστικών μονάδων. Το πρόβλημα είναι ότι υπάρχουν προφανείς περιορισμοί στην παράλληλη εκτέλεση μιας διαδοχικής ροής εντολών, επομένως η τυφλή αύξηση του αριθμού των υπολογιστικών μονάδων δεν παρέχει κανένα όφελος, καθώς θα εξακολουθούν να είναι αδρανείς τις περισσότερες φορές.

Αντίθετα, η λειτουργία της GPU είναι σχετικά απλή. Αποτελείται από τη λήψη μιας ομάδας πολυγώνων στη μία πλευρά και τη δημιουργία μιας ομάδας εικονοστοιχείων από την άλλη. Τα πολύγωνα και τα εικονοστοιχεία είναι ανεξάρτητα μεταξύ τους, επομένως μπορούν να υποβληθούν σε παράλληλη επεξεργασία. Έτσι, σε μια GPU είναι δυνατό να εκχωρηθεί ένα μεγάλο μέρος του κρυστάλλου σε υπολογιστικές μονάδες, οι οποίες, σε αντίθεση με την CPU, θα χρησιμοποιηθούν στην πραγματικότητα.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Η GPU διαφέρει από την CPU όχι μόνο με αυτόν τον τρόπο. Η πρόσβαση στη μνήμη στη GPU είναι πολύ συνδεδεμένη - εάν διαβάζεται ένα texel, τότε μετά από μερικούς κύκλους ρολογιού θα διαβάζεται το γειτονικό texel. Όταν εγγράφεται ένα pixel, μετά από μερικούς κύκλους ρολογιού θα εγγραφεί το γειτονικό. Οργανώνοντας έξυπνα τη μνήμη, μπορείτε να επιτύχετε απόδοση κοντά στη θεωρητική απόδοση. Αυτό σημαίνει ότι η GPU, σε αντίθεση με την CPU, δεν απαιτεί τεράστια κρυφή μνήμη, αφού ο ρόλος της είναι να επιταχύνει τις λειτουργίες δημιουργίας υφής. Το μόνο που χρειάζεται είναι μερικά kilobyte που περιέχουν μερικά texel που χρησιμοποιούνται σε διγραμμικά και τριγραμμικά φίλτρα.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Ζήτω το GeForce FX!

Οι δύο κόσμοι έμειναν χωρισμένοι για πολύ καιρό. Χρησιμοποιήσαμε CPU (ή ακόμα και πολλαπλές CPU) για εργασίες γραφείου και εφαρμογές Διαδικτύου και οι GPU ήταν καλές μόνο για την επιτάχυνση της απόδοσης. Αλλά ένα χαρακτηριστικό άλλαξε τα πάντα: δηλαδή, η έλευση των προγραμματιζόμενων GPU. Στην αρχή, οι CPU δεν είχαν τίποτα να φοβηθούν. Οι πρώτες λεγόμενες προγραμματιζόμενες GPU (NV20 και R200) ​​δεν αποτελούσαν σχεδόν απειλή. Ο αριθμός των εντολών στο πρόγραμμα παρέμεινε περιορισμένος σε περίπου 10 και δούλευαν σε πολύ εξωτικούς τύπους δεδομένων, όπως αριθμούς σταθερού σημείου 9 ή 12 bit.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Αλλά ο νόμος του Μουρ έδειξε και πάλι την καλύτερή του πλευρά. Η αύξηση του αριθμού των τρανζίστορ όχι μόνο κατέστησε δυνατή την αύξηση του αριθμού των υπολογιστικών μονάδων, αλλά και βελτίωσε την ευελιξία τους. Η εμφάνιση του NV30 μπορεί να θεωρηθεί σημαντικό βήμα προόδου για διάφορους λόγους. Φυσικά, στους παίκτες δεν άρεσαν πολύ οι κάρτες NV30, αλλά οι νέες GPU άρχισαν να βασίζονται σε δύο χαρακτηριστικά που είχαν σχεδιαστεί για να αλλάξουν την αντίληψη των GPU ως κάτι περισσότερο από απλά επιταχυντές γραφικών.

• Υποστήριξη για υπολογισμούς κινητής υποδιαστολής μονής ακρίβειας (ακόμα και αν δεν συμμορφωνόταν με το πρότυπο IEEE754).
• υποστήριξη για περισσότερες από χίλιες οδηγίες.

Έχουμε λοιπόν όλες τις προϋποθέσεις που μπορούν να προσελκύσουν πρωτοπόρους ερευνητές που αναζητούν πάντα πρόσθετη υπολογιστική ισχύ.

Η ιδέα της χρήσης επιταχυντών γραφικών για μαθηματικούς υπολογισμούς δεν είναι νέα. Οι πρώτες προσπάθειες έγιναν στη δεκαετία του '90 του περασμένου αιώνα. Φυσικά, ήταν πολύ πρωτόγονοι - περιορίζονταν, ως επί το πλείστον, στη χρήση ορισμένων λειτουργιών υλικού, όπως η ραστεροποίηση και τα Z-buffers, για την επιτάχυνση εργασιών όπως η εύρεση διαδρομής ή το συμπέρασμα Διαγράμματα Voronoi .

Κάντε κλικ στην εικόνα για μεγέθυνση.

Το 2003, με την εμφάνιση των εξελιγμένων shaders, δημιουργήθηκε μια νέα γραμμή - αυτή τη φορά εκτελώντας υπολογισμούς μήτρας. Αυτή ήταν η χρονιά που μια ολόκληρη ενότητα του SIGGRAPH ("Υπολογισμοί στις GPU") αφιερώθηκε σε μια νέα περιοχή της πληροφορικής. Αυτή η πρώιμη πρωτοβουλία ονομάστηκε GPGPU (υπολογισμός γενικού σκοπού σε GPU). Και ένα πρώιμο σημείο καμπής ήταν η εμφάνιση του .

Για να κατανοήσετε τον ρόλο του BrookGPU, πρέπει να καταλάβετε πώς συνέβησαν όλα πριν από την εμφάνισή του. Ο μόνος τρόπος για να αποκτήσετε πόρους GPU το 2003 ήταν να χρησιμοποιήσετε ένα από τα δύο API γραφικών - Direct3D ή OpenGL. Κατά συνέπεια, οι προγραμματιστές που ήθελαν δυνατότητες GPU για τον υπολογισμό τους έπρεπε να βασίζονται στα δύο αναφερόμενα API. Το πρόβλημα είναι ότι δεν ήταν πάντα ειδικοί στον προγραμματισμό καρτών γραφικών, και αυτό δυσκόλεψε σοβαρά την πρόσβαση στην τεχνολογία. Εάν οι προγραμματιστές 3D λειτουργούν με shaders, υφές και θραύσματα, τότε οι ειδικοί στον τομέα του παράλληλου προγραμματισμού βασίζονται σε νήματα, πυρήνες, scatter κ.λπ. Ως εκ τούτου, πρώτα ήταν απαραίτητο να γίνουν αναλογίες μεταξύ των δύο κόσμων.

• Ρεύμαείναι μια ροή στοιχείων του ίδιου τύπου στη GPU που μπορεί να αναπαρασταθεί από μια υφή. Κατ 'αρχήν, στον κλασικό προγραμματισμό υπάρχει ένα τέτοιο ανάλογο ως πίνακας.
• Πυρήνας- μια λειτουργία που θα εφαρμοστεί ανεξάρτητα σε κάθε στοιχείο του ρεύματος. είναι το ισοδύναμο ενός pixel shader. Στον κλασικό προγραμματισμό, μπορούμε να δώσουμε μια αναλογία ενός βρόχου - εφαρμόζεται σε μεγάλο αριθμό στοιχείων.
• Για να διαβάσετε τα αποτελέσματα της εφαρμογής του πυρήνα σε ένα νήμα, πρέπει να δημιουργηθεί μια υφή. Δεν υπάρχει αντίστοιχο στην CPU, αφού έχει πρόσβαση σε πλήρη μνήμη.
• Ο έλεγχος της θέσης στη μνήμη όπου θα γίνει η εγγραφή (σε πράξεις διασποράς) γίνεται μέσω του σκίαστρου κορυφής, αφού ο σκιαστής εικονοστοιχείων δεν μπορεί να αλλάξει τις συντεταγμένες του υπό επεξεργασία εικονοστοιχείου.

Όπως μπορείτε να δείτε, ακόμη και αν ληφθούν υπόψη οι παραπάνω αναλογίες, η εργασία δεν φαίνεται απλή. Και ο Μπρουκ ήρθε στη διάσωση. Αυτό το όνομα αναφέρεται σε επεκτάσεις στη γλώσσα C ("C με ροές", "C με ροές"), όπως τις ονόμασαν οι προγραμματιστές στο Stanford. Στον πυρήνα του, το καθήκον του Brook ήταν να κρύψει όλα τα στοιχεία του 3D API από τον προγραμματιστή, γεγονός που επέτρεψε την παρουσίαση της GPU ως συνεπεξεργαστή για παράλληλους υπολογιστές. Για να γίνει αυτό, ο μεταγλωττιστής Brook επεξεργάστηκε ένα αρχείο .br με κώδικα και επεκτάσεις C++ και στη συνέχεια δημιούργησε κώδικα C++ που συνδέθηκε σε μια βιβλιοθήκη με υποστήριξη για διαφορετικές εξόδους (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Κάντε κλικ στην εικόνα για μεγέθυνση.

Ο Brook έχει αρκετές πιστώσεις στο ενεργητικό του, ο πρώτος από τους οποίους είναι να φέρει το GPGPU από τη σκιά, ώστε η τεχνολογία να μπορεί να αγκαλιαστεί από τις μάζες. Αν και μετά την ανακοίνωση του έργου, μια σειρά από ιστότοπους πληροφορικής ανέφεραν υπερβολικά αισιόδοξα ότι η κυκλοφορία του Brook θέτει αμφιβολίες για την ύπαρξη CPU, οι οποίες σύντομα θα αντικατασταθούν από πιο ισχυρές GPU. Όμως, όπως βλέπουμε, ακόμη και μετά από πέντε χρόνια αυτό δεν συνέβη. Ειλικρινά, δεν πιστεύουμε ότι αυτό θα συμβεί ποτέ. Από την άλλη, εξετάζοντας την επιτυχημένη εξέλιξη των CPU, οι οποίες προσανατολίζονται όλο και περισσότερο στον παραλληλισμό (περισσότεροι πυρήνες, τεχνολογία πολλαπλών νημάτων SMT, επέκταση μπλοκ SIMD), καθώς και των GPU, οι οποίες, αντίθετα, γίνονται όλο και πιο καθολικές (υποστήριξη για υπολογισμούς κινητής υποδιαστολής) απλή ακρίβεια, υπολογισμοί ακέραιων αριθμών, υποστήριξη υπολογισμών διπλής ακρίβειας), φαίνεται ότι η GPU και η CPU σύντομα απλά θα συγχωνευθούν. Τι θα γίνει τότε; Θα απορροφηθούν οι GPU από τις CPU, όπως συνέβη με τους μαθηματικούς συνεπεξεργαστές; Πολύ πιθανό. Η Intel και η AMD εργάζονται σε παρόμοια έργα σήμερα. Αλλά πολλά μπορούν ακόμα να αλλάξουν.

Ας επιστρέψουμε όμως στο θέμα μας. Το πλεονέκτημα του Brook ήταν να διαδώσει την έννοια του GPGPU, απλοποίησε σημαντικά την πρόσβαση σε πόρους GPU, γεγονός που επέτρεψε σε όλο και περισσότερους χρήστες να κυριαρχήσουν στο νέο μοντέλο προγραμματισμού. Από την άλλη πλευρά, παρά όλες τις ιδιότητες του Brook, υπήρχε ακόμη πολύς δρόμος για να χρησιμοποιηθούν οι πόροι της GPU για υπολογιστές.

Ένα από τα προβλήματα σχετίζεται με τα διαφορετικά επίπεδα αφαίρεσης, και επίσης, ειδικότερα, με το υπερβολικό πρόσθετο φορτίο που δημιουργείται από το 3D API, το οποίο μπορεί να είναι αρκετά αισθητό. Αλλά ένα πιο σοβαρό πρόβλημα μπορεί να θεωρηθεί πρόβλημα συμβατότητας, με το οποίο οι προγραμματιστές του Brook δεν μπορούσαν να κάνουν τίποτα. Υπάρχει έντονος ανταγωνισμός μεταξύ των κατασκευαστών GPU, επομένως συχνά βελτιστοποιούν τους οδηγούς τους. Ενώ βελτιστοποιήσεις όπως αυτές είναι κυρίως καλές για τους παίκτες, θα μπορούσαν να τερματίσουν τη συμβατότητα του Brook εν μία νυκτί. Επομένως, είναι δύσκολο να φανταστεί κανείς τη χρήση αυτού του API σε κώδικα παραγωγής που θα λειτουργήσει κάπου. Και για πολύ καιρό, ο Μπρουκ παρέμεινε κτήμα ερασιτεχνών ερευνητών και προγραμματιστών.

Ωστόσο, η επιτυχία του Brook ήταν αρκετή για να τραβήξει την προσοχή της ATI και της nVidia, που ενδιαφέρθηκαν για μια τέτοια πρωτοβουλία, καθώς θα μπορούσε να επεκτείνει την αγορά, ανοίγοντας έναν νέο σημαντικό τομέα για τις εταιρείες.

Οι ερευνητές που συμμετείχαν αρχικά στο έργο Brook γρήγορα ενώθηκαν με ομάδες ανάπτυξης στη Σάντα Κλάρα για να παρουσιάσουν μια παγκόσμια στρατηγική για την ανάπτυξη της νέας αγοράς. Η ιδέα ήταν να δημιουργηθεί ένας συνδυασμός υλικού και λογισμικού κατάλληλου για εργασίες GPGPU. Δεδομένου ότι οι προγραμματιστές της nVidia γνωρίζουν όλα τα μυστικά των GPU τους, δεν μπορούσαν να βασιστούν στο API γραφικών, αλλά να επικοινωνήσουν με τη GPU μέσω ενός προγράμματος οδήγησης. Αν και, φυσικά, αυτό έρχεται με τα δικά του προβλήματα. Έτσι, η ομάδα ανάπτυξης CUDA (Compute Unified Device Architecture) δημιούργησε ένα σύνολο επιπέδων λογισμικού για εργασία με GPU.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Όπως μπορείτε να δείτε στο διάγραμμα, το CUDA παρέχει δύο API.

• API υψηλού επιπέδου: CUDA Runtime API.
• χαμηλού επιπέδου API: CUDA Driver API.

Επειδή το API υψηλού επιπέδου υλοποιείται πάνω από το API χαμηλού επιπέδου, κάθε κλήση σε μια συνάρτηση σε επίπεδο χρόνου εκτέλεσης αναλύεται σε απλούστερες οδηγίες που επεξεργάζεται το API προγράμματος οδήγησης. Λάβετε υπόψη ότι τα δύο API είναι αμοιβαία αποκλειόμενα: ένας προγραμματιστής μπορεί να χρησιμοποιήσει το ένα ή το άλλο API, αλλά δεν θα είναι δυνατός ο συνδυασμός κλήσεων συναρτήσεων από τα δύο API. Γενικά, ο όρος "API υψηλού επιπέδου" είναι σχετικός. Ακόμη και το Runtime API είναι τέτοιο που πολλοί θα το θεωρούσαν χαμηλού επιπέδου. Ωστόσο, εξακολουθεί να παρέχει λειτουργίες που είναι πολύ βολικές για την προετοιμασία ή τη διαχείριση του περιβάλλοντος. Αλλά μην περιμένετε ένα ιδιαίτερα υψηλό επίπεδο αφαίρεσης - πρέπει να έχετε ακόμα μια καλή γνώση σχετικά με τις GPU της nVidia και τον τρόπο λειτουργίας τους.

Το API του προγράμματος οδήγησης είναι ακόμα πιο δύσκολο να δουλέψεις. θα χρειαστεί περισσότερη προσπάθεια για την εκτέλεση της επεξεργασίας GPU. Από την άλλη πλευρά, ένα API χαμηλού επιπέδου είναι πιο ευέλικτο, δίνοντας στον προγραμματιστή πρόσθετο έλεγχο εάν χρειάζεται. Δύο API μπορούν να λειτουργούν με πόρους OpenGL ή Direct3D (μόνο η ένατη έκδοση από σήμερα). Το όφελος αυτής της δυνατότητας είναι προφανές - το CUDA μπορεί να χρησιμοποιηθεί για τη δημιουργία πόρων (γεωμετρία, διαδικαστικές υφές, κ.λπ.) που μπορούν να περάσουν στο API γραφικών ή, αντίθετα, μπορείτε να ζητήσετε από το 3D API να στείλει αποτελέσματα απόδοσης στο πρόγραμμα CUDA, το οποίο, με τη σειρά του, θα εκτελέσει μετα-επεξεργασία. Υπάρχουν πολλά παραδείγματα τέτοιων αλληλεπιδράσεων και το πλεονέκτημα είναι ότι οι πόροι συνεχίζουν να αποθηκεύονται στη μνήμη της GPU, δεν χρειάζεται να μεταφερθούν μέσω του διαύλου PCI Express, ο οποίος εξακολουθεί να παραμένει εμπόδιο.

Ωστόσο, πρέπει να σημειωθεί ότι η κοινή χρήση πόρων στη μνήμη βίντεο δεν είναι πάντα ιδανική και μπορεί να οδηγήσει σε κάποιους πονοκεφάλους. Για παράδειγμα, κατά την αλλαγή της ανάλυσης ή του βάθους χρώματος, τα δεδομένα γραφικών έχουν προτεραιότητα. Επομένως, εάν χρειαστεί να αυξήσετε τους πόρους στην προσωρινή μνήμη πλαισίων, το πρόγραμμα οδήγησης θα το κάνει εύκολα εις βάρος των πόρων των εφαρμογών CUDA, οι οποίες απλώς θα κολλήσουν με ένα σφάλμα. Φυσικά, όχι πολύ κομψό, αλλά μια τέτοια κατάσταση δεν πρέπει να συμβαίνει πολύ συχνά. Και αφού αρχίσαμε να μιλάμε για μειονεκτήματα: εάν θέλετε να χρησιμοποιήσετε πολλές GPU για εφαρμογές CUDA, τότε πρέπει πρώτα να απενεργοποιήσετε τη λειτουργία SLI, διαφορετικά οι εφαρμογές CUDA θα μπορούν να «βλέπουν» μόνο μία GPU.

Τέλος, το τρίτο επίπεδο λογισμικού είναι αφιερωμένο στις βιβλιοθήκες - δύο, για την ακρίβεια.

• CUBLAS, το οποίο περιέχει τα απαραίτητα μπλοκ για υπολογισμούς γραμμικής άλγεβρας στη GPU.
• CUFFT, το οποίο υποστηρίζει τον υπολογισμό των μετασχηματισμών Fourier - ένας αλγόριθμος που χρησιμοποιείται ευρέως στον τομέα της επεξεργασίας σήματος.

Πριν βουτήξουμε στο CUDA, ας ορίσουμε έναν αριθμό όρων διάσπαρτων στην τεκμηρίωση της nVidia. Η εταιρεία έχει επιλέξει μια πολύ συγκεκριμένη ορολογία που δύσκολα συνηθίζεις. Πρώτα από όλα, σημειώνουμε ότι Νήμαστο CUDA δεν έχει την ίδια σημασία με ένα νήμα CPU, και επίσης δεν είναι το ισοδύναμο ενός νήματος στα άρθρα μας για τις GPU. Το νήμα της GPU σε αυτήν την περίπτωση είναι το βασικό σύνολο δεδομένων που πρέπει να υποβληθούν σε επεξεργασία. Σε αντίθεση με τα νήματα της CPU, τα νήματα CUDA είναι πολύ «ελαφριά», δηλαδή, η εναλλαγή περιβάλλοντος μεταξύ δύο νημάτων δεν είναι μια λειτουργία έντασης πόρων.

Ο δεύτερος όρος που συναντάται συχνά στην τεκμηρίωση CUDA είναι στημόνι. Δεν υπάρχει καμία σύγχυση εδώ, αφού δεν υπάρχει ανάλογο στα ρωσικά (εκτός αν είστε λάτρης του Start Trek ή του παιχνιδιού Warhammer). Στην πραγματικότητα, ο όρος προέρχεται από την κλωστοϋφαντουργία, όπου το νήμα υφαδιού τραβιέται μέσω ενός νήματος στημονιού που τεντώνεται σε έναν αργαλειό. Το warp στο CUDA είναι μια ομάδα 32 νημάτων και είναι η ελάχιστη ποσότητα δεδομένων που υποβάλλονται σε επεξεργασία με τρόπο SIMD σε πολυεπεξεργαστές CUDA.

Αλλά μια τέτοια "κοκκία" δεν είναι πάντα βολική για τον προγραμματιστή. Επομένως, στο CUDA, αντί να εργάζεστε απευθείας με στημόνι, μπορείτε να εργαστείτε με μπλοκ, που περιέχει από 64 έως 512 νήματα.

Τέλος, αυτά τα μπλοκ ενώνονται πλέγμα. Το πλεονέκτημα αυτής της ομαδοποίησης είναι ότι ο αριθμός των μπλοκ που επεξεργάζεται ταυτόχρονα η GPU σχετίζεται στενά με τους πόρους υλικού, όπως θα δούμε παρακάτω. Η ομαδοποίηση μπλοκ σε πλέγματα σάς επιτρέπει να αφαιρέσετε εντελώς αυτόν τον περιορισμό και να εφαρμόσετε τον πυρήνα σε περισσότερα νήματα σε μία κλήση χωρίς να ανησυχείτε για σταθερούς πόρους. Υπεύθυνες για όλα αυτά είναι οι βιβλιοθήκες CUDA. Επιπλέον, ένα τέτοιο μοντέλο κλιμακώνεται καλά. Εάν η GPU έχει λίγους πόρους, θα εκτελεί μπλοκ διαδοχικά. Εάν ο αριθμός των υπολογιστικών επεξεργαστών είναι μεγάλος, τότε τα μπλοκ μπορούν να εκτελεστούν παράλληλα. Δηλαδή, ο ίδιος κώδικας μπορεί να εκτελεστεί σε αρχικές GPU, καθώς και σε κορυφαία και ακόμη και μελλοντικά μοντέλα.

Υπάρχουν μερικοί ακόμη όροι στο CUDA API που υποδεικνύουν CPU ( οικοδεσπότης/οικοδεσπότης) και GPU ( συσκευή). Εάν αυτή η μικρή εισαγωγή δεν σας έχει τρομάξει, τότε ήρθε η ώρα να ρίξετε μια πιο προσεκτική ματιά στο CUDA.

Εάν διαβάζετε τακτικά τον Οδηγό υλικού του Tom, τότε η αρχιτεκτονική των πιο πρόσφατων GPU της nVidia είναι οικεία σε εσάς. Εάν όχι, σας συνιστούμε να διαβάσετε το άρθρο ". nVidia GeForce GTX 260 και 280: νέα γενιά καρτών γραφικών«Όταν πρόκειται για το CUDA, η Nvidia παρουσιάζει την αρχιτεκτονική λίγο διαφορετικά, αποκαλύπτοντας κάποιες λεπτομέρειες που προηγουμένως ήταν κρυμμένες.

Όπως μπορείτε να δείτε από την παραπάνω εικόνα, ο πυρήνας shader nVidia αποτελείται από πολλά συμπλέγματα επεξεργαστών υφής (Σύμπλεγμα επεξεργαστή υφής, TPC). Η κάρτα γραφικών 8800 GTX, για παράδειγμα, χρησιμοποιούσε οκτώ συμπλέγματα, η 8800 GTS - έξι κ.λπ. Κάθε σύμπλεγμα αποτελείται ουσιαστικά από ένα μπλοκ υφής και δύο πολυεπεξεργαστή ροής. Οι τελευταίες περιλαμβάνουν την αρχή του αγωγού (μπροστινό άκρο), που διαβάζει και αποκωδικοποιεί οδηγίες, καθώς και την αποστολή τους για εκτέλεση, και το τέλος του αγωγού (πίσω άκρο), που αποτελείται από οκτώ υπολογιστικές συσκευές και δύο υπερλειτουργικές συσκευές SFU (Super Function Unit), όπου οι εντολές εκτελούνται χρησιμοποιώντας την αρχή SIMD, δηλαδή μια εντολή εφαρμόζεται σε όλα τα νήματα στο στημόνι. Η nVidia καλεί αυτήν τη μέθοδο εκτέλεσης SIMT(μία εντολή πολλαπλά νήματα, μία εντολή, πολλά νήματα). Είναι σημαντικό να σημειωθεί ότι το άκρο του μεταφορέα λειτουργεί με διπλάσια συχνότητα από την αρχή του. Στην πράξη, αυτό σημαίνει ότι το εξάρτημα εμφανίζεται δύο φορές πιο «πλατύ» από ότι είναι στην πραγματικότητα (δηλαδή ως μπλοκ SIMD 16 καναλιών αντί για οκτώ καναλιών). Οι πολυεπεξεργαστές ροής λειτουργούν ως εξής: κάθε κύκλος ρολογιού, η αρχή του αγωγού επιλέγει ένα warp έτοιμο για εκτέλεση και ξεκινά την εκτέλεση της εντολής. Για να εφαρμοστεί η οδηγία και στα 32 νήματα στο στημόνι, το τέλος του αγωγού θα απαιτήσει τέσσερις κύκλους ρολογιού, αλλά δεδομένου ότι τρέχει με διπλάσια συχνότητα από την αρχή, θα απαιτήσει μόνο δύο κύκλους ρολογιού (όσον αφορά την αρχή του ο αγωγός). Επομένως, ώστε η αρχή του αγωγού να μην αδράνει έναν κύκλο ρολογιού και το υλικό να φορτώνεται στο μέγιστο, σε ιδανική περίπτωση, μπορείτε να εναλλάσσετε οδηγίες σε κάθε κύκλο ρολογιού - μια κλασική οδηγία σε έναν κύκλο ρολογιού και μια οδηγία για SFU σε έναν άλλο .

Κάθε πολυεπεξεργαστής έχει ένα συγκεκριμένο σύνολο πόρων που αξίζει να κατανοήσουμε. Υπάρχει μια μικρή περιοχή μνήμης που ονομάζεται "Κοινή μνήμη", 16 kbyte ανά πολυεπεξεργαστή. Αυτή δεν είναι καθόλου προσωρινή μνήμη: ο προγραμματιστής μπορεί να τη χρησιμοποιήσει κατά την κρίση του. Δηλαδή, έχουμε κάτι κοντά στο Local Store of SPUs σε επεξεργαστές Cell. Αυτή η λεπτομέρεια είναι πολύ ενδιαφέρουσα, αφού τονίζει ότι το CUDA είναι ένας συνδυασμός τεχνολογιών λογισμικού και υλικού. Αυτή η περιοχή μνήμης δεν χρησιμοποιείται για σκίαση εικονοστοιχείων, κάτι που η Nvidia επισημαίνει έξυπνα «δεν μας αρέσει τα εικονοστοιχεία να μιλάνε μεταξύ τους».

Αυτή η περιοχή μνήμης ανοίγει τη δυνατότητα ανταλλαγής πληροφοριών μεταξύ νημάτων σε ένα μπλοκ. Είναι σημαντικό να τονίσουμε αυτόν τον περιορισμό: όλα τα νήματα σε ένα μπλοκ είναι εγγυημένα ότι θα εκτελεστούν από έναν μόνο πολυεπεξεργαστή. Αντίθετα, η εκχώρηση μπλοκ σε διαφορετικούς πολυεπεξεργαστές δεν καθορίζεται καθόλου και δύο νήματα από διαφορετικά μπλοκ δεν μπορούν να ανταλλάξουν πληροφορίες μεταξύ τους κατά την εκτέλεση. Δηλαδή, η χρήση κοινής μνήμης δεν είναι τόσο εύκολη. Ωστόσο, η κοινόχρηστη μνήμη εξακολουθεί να δικαιολογείται, εκτός από τις περιπτώσεις όπου πολλά νήματα προσπαθούν να αποκτήσουν πρόσβαση στην ίδια τράπεζα μνήμης, προκαλώντας διένεξη. Σε άλλες περιπτώσεις, η πρόσβαση στην κοινόχρηστη μνήμη είναι εξίσου γρήγορη με την πρόσβαση σε καταχωρητές.

Η κοινόχρηστη μνήμη δεν είναι η μόνη μνήμη στην οποία μπορούν να έχουν πρόσβαση οι πολυεπεξεργαστές. Μπορούν να χρησιμοποιούν μνήμη βίντεο, αλλά με μικρότερο εύρος ζώνης και υψηλότερο λανθάνοντα χρόνο. Επομένως, προκειμένου να μειωθεί η συχνότητα πρόσβασης σε αυτή τη μνήμη, η nVidia εξόπλισε πολυεπεξεργαστές με μια κρυφή μνήμη (περίπου 8 KB ανά πολυεπεξεργαστή) που αποθηκεύει σταθερές και υφές.

Ο πολυεπεξεργαστής έχει 8.192 καταχωρητές που είναι κοινοί σε όλα τα νήματα όλων των μπλοκ που είναι ενεργά στον πολυεπεξεργαστή. Ο αριθμός των ενεργών μπλοκ ανά πολυεπεξεργαστή δεν μπορεί να υπερβαίνει τα οκτώ και ο αριθμός των ενεργών παραμορφώσεων περιορίζεται σε 24 (768 νήματα). Επομένως, η 8800 GTX μπορεί να χειριστεί έως και 12.288 νήματα τη φορά. Όλοι αυτοί οι περιορισμοί αξίζει να αναφερθούν γιατί επιτρέπουν τη βελτιστοποίηση του αλγόριθμου με βάση τους διαθέσιμους πόρους.

Η βελτιστοποίηση ενός προγράμματος CUDA συνίσταται επομένως στην επίτευξη μιας βέλτιστης ισορροπίας μεταξύ του αριθμού των μπλοκ και του μεγέθους τους. Περισσότερα νήματα ανά μπλοκ θα είναι χρήσιμα για τη μείωση του λανθάνοντος χρόνου της μνήμης, αλλά μειώνεται επίσης ο αριθμός των διαθέσιμων καταχωρητών ανά νήμα. Επιπλέον, ένα μπλοκ 512 νημάτων θα είναι αναποτελεσματικό επειδή μόνο ένα μπλοκ μπορεί να είναι ενεργό σε έναν πολυεπεξεργαστή, με αποτέλεσμα την απώλεια 256 νημάτων. Ως εκ τούτου, η nVidia συνιστά τη χρήση μπλοκ 128 ή 256 νημάτων, γεγονός που παρέχει τον βέλτιστο συμβιβασμό μεταξύ μειωμένου λανθάνοντος χρόνου και του αριθμού των καταχωρητών για τους περισσότερους πυρήνες/πυρήνες.

Από την άποψη του λογισμικού, το CUDA αποτελείται από ένα σύνολο επεκτάσεων στη γλώσσα C, που θυμίζει BrookGPU, καθώς και από πολλές συγκεκριμένες κλήσεις API. Μεταξύ των επεκτάσεων υπάρχουν προσδιοριστές τύπου που σχετίζονται με συναρτήσεις και μεταβλητές. Είναι σημαντικό να θυμάστε τη λέξη-κλειδί __παγκόσμια__, η οποία, όταν δίνεται πριν από τη συνάρτηση, δείχνει ότι η τελευταία ανήκει στον πυρήνα - αυτή η συνάρτηση θα κληθεί από την CPU και θα εκτελεστεί στη GPU. Πρόθεμα __συσκευή__υποδεικνύει ότι η συνάρτηση θα εκτελεστεί στη GPU (η οποία, παρεμπιπτόντως, είναι αυτό που το CUDA αποκαλεί "συσκευή") αλλά μπορεί να κληθεί μόνο από τη GPU (με άλλα λόγια, από μια άλλη λειτουργία __device__ ή από μια συνάρτηση __global__) . Τέλος, το πρόθεμα __πλήθος__προαιρετικό, υποδηλώνει μια συνάρτηση που καλείται από την CPU και εκτελείται από την CPU - με άλλα λόγια, μια κανονική συνάρτηση.

Υπάρχει ένας αριθμός περιορισμών που σχετίζονται με τις συναρτήσεις __device__ και __global__: δεν μπορούν να είναι αναδρομικές (δηλαδή να καλούν τον εαυτό τους) και δεν μπορούν να έχουν μεταβλητό αριθμό ορισμάτων. Τέλος, δεδομένου ότι οι λειτουργίες __device__ βρίσκονται στο χώρο της μνήμης GPU, είναι λογικό να μην μπορεί να ληφθεί η διεύθυνσή τους. Οι μεταβλητές έχουν επίσης έναν αριθμό χαρακτηριστικών που υποδεικνύουν την περιοχή μνήμης όπου θα αποθηκευτούν. Μεταβλητή με πρόθεμα __κοινόχρηστο__σημαίνει ότι θα αποθηκευτεί στην κοινόχρηστη μνήμη του πολυεπεξεργαστή ροής. Η κλήση συνάρτησης __global__ είναι ελαφρώς διαφορετική. Το θέμα είναι ότι κατά την κλήση, πρέπει να ορίσετε τη διαμόρφωση εκτέλεσης - πιο συγκεκριμένα, το μέγεθος του πλέγματος στο οποίο θα εφαρμοστεί ο πυρήνας, καθώς και το μέγεθος κάθε μπλοκ. Ας πάρουμε, για παράδειγμα, έναν πυρήνα με την ακόλουθη υπογραφή.

__global__ void Func(παράμετρος float*);

Θα ονομάζεται ως

Func<<< Dg, Db >>>(παράμετρος);

όπου Dg είναι το μέγεθος πλέγματος και Db το μέγεθος του μπλοκ. Αυτές οι δύο μεταβλητές αναφέρονται σε έναν νέο τύπο διανύσματος που εισήχθη με το CUDA.

Το CUDA API περιέχει λειτουργίες για εργασία με μνήμη σε VRAM: cudaMalloc για εκχώρηση μνήμης, cudaFree για την απελευθέρωσή της και cudaMemcpy για αντιγραφή μνήμης μεταξύ RAM και VRAM και αντίστροφα.

Θα ολοκληρώσουμε αυτήν την ανασκόπηση με έναν πολύ ενδιαφέροντα τρόπο με τον οποίο συντάσσεται ένα πρόγραμμα CUDA: η μεταγλώττιση πραγματοποιείται σε διάφορα στάδια. Αρχικά, εξάγεται ο ειδικός κώδικας της CPU και μεταβιβάζεται στον τυπικό μεταγλωττιστή. Ο κώδικας που προορίζεται για την GPU μετατρέπεται πρώτα στην ενδιάμεση γλώσσα PTX. Είναι παρόμοιο με ένα assembler και σας επιτρέπει να εξετάσετε τον κώδικα αναζητώντας πιθανές ανεπάρκειες. Τέλος, η τελευταία φάση είναι η μετάφραση της ενδιάμεσης γλώσσας σε συγκεκριμένες οδηγίες GPU και η δημιουργία ενός δυαδικού αρχείου.

Αφού κοίταξα την τεκμηρίωση της nVidia, μπαίνω στον πειρασμό να δοκιμάσω το CUDA αυτή την εβδομάδα. Αλήθεια, τι καλύτερο από την αξιολόγηση ενός API δημιουργώντας το δικό σας πρόγραμμα; Αυτή είναι η στιγμή που τα περισσότερα προβλήματα θα πρέπει να εμφανιστούν, ακόμα κι αν όλα φαίνονται τέλεια στο χαρτί. Επίσης, η πρακτική θα δείξει καλύτερα πόσο καλά έχετε κατανοήσει όλες τις αρχές που περιγράφονται στην τεκμηρίωση CUDA.

Είναι πολύ εύκολο να εμπλακείς σε ένα έργο σαν αυτό. Σήμερα υπάρχει ένας μεγάλος αριθμός δωρεάν αλλά υψηλής ποιότητας εργαλείων διαθέσιμα για λήψη. Για τη δοκιμή μας χρησιμοποιήσαμε το Visual C++ Express 2005, το οποίο έχει όλα όσα χρειαζόμαστε. Το πιο δύσκολο κομμάτι ήταν να βρούμε ένα πρόγραμμα που δεν θα χρειαζόταν εβδομάδες για να μεταφερθεί στη GPU, αλλά θα ήταν επίσης αρκετά ενδιαφέρον ώστε οι προσπάθειές μας να μην ήταν μάταιες. Στο τέλος, επιλέξαμε ένα κομμάτι κώδικα που παίρνει έναν χάρτη ύψους και υπολογίζει έναν αντίστοιχο κανονικό χάρτη. Δεν θα υπεισέλθουμε σε λεπτομέρειες σχετικά με αυτήν τη λειτουργία, καθώς δεν είναι καθόλου ενδιαφέρουσα σε αυτό το άρθρο. Για να είμαστε συνοπτικά, το πρόγραμμα ασχολείται με την καμπυλότητα των περιοχών: για κάθε εικονοστοιχείο της αρχικής εικόνας, τοποθετούμε έναν πίνακα που καθορίζει το χρώμα του προκύπτοντος εικονοστοιχείου στην εικόνα που δημιουργείται από γειτονικά εικονοστοιχεία, χρησιμοποιώντας έναν περισσότερο ή λιγότερο περίπλοκο τύπο. Το πλεονέκτημα αυτής της λειτουργίας είναι ότι είναι πολύ εύκολο να παραλληλιστεί, επομένως αυτή η δοκιμή δείχνει τέλεια τις δυνατότητες του CUDA.

Ένα άλλο πλεονέκτημα είναι ότι έχουμε ήδη μια υλοποίηση CPU, οπότε μπορούμε να συγκρίνουμε το αποτέλεσμα με την έκδοση CUDA - χωρίς να εφεύρουμε ξανά τον τροχό.

Ας επαναλάβουμε για άλλη μια φορά ότι σκοπός της δοκιμής ήταν η γνωριμία με τα βοηθητικά προγράμματα CUDA SDK και όχι η σύγκριση των εκδόσεων για CPU και GPU. Δεδομένου ότι αυτή ήταν η πρώτη μας προσπάθεια να δημιουργήσουμε ένα πρόγραμμα CUDA, δεν περιμέναμε πραγματικά υψηλή απόδοση. Δεδομένου ότι αυτό το μέρος του κώδικα δεν είναι κρίσιμο, η έκδοση για τη CPU δεν βελτιστοποιήθηκε, επομένως μια άμεση σύγκριση των αποτελεσμάτων δεν είναι καθόλου ενδιαφέρουσα.

Εκτέλεση

Ωστόσο, μετρήσαμε τον χρόνο εκτέλεσης για να δούμε αν υπάρχει πλεονέκτημα στη χρήση του CUDA ακόμη και με την πιο ωμή υλοποίηση ή αν χρειαζόμασταν μακρά και κουραστική πρακτική για να έχουμε οποιοδήποτε όφελος κατά τη χρήση της GPU. Το μηχάνημα δοκιμής ελήφθη από το εργαστήριο ανάπτυξης μας - ένας φορητός υπολογιστής με επεξεργαστή Core 2 Duo T5450 και μια κάρτα γραφικών GeForce 8600M GT με Vista. Αυτό απέχει πολύ από έναν υπερυπολογιστή, αλλά τα αποτελέσματα είναι πολύ ενδιαφέροντα, καθώς η δοκιμή δεν είναι «προσαρμοσμένη» για τη GPU. Είναι πάντα ωραίο να βλέπεις τη Nvidia να παρουσιάζει τεράστια κέρδη σε συστήματα με τερατώδεις GPU και άφθονο εύρος ζώνης, αλλά στην πράξη πολλές από τις 70 εκατομμύρια GPU με δυνατότητα CUDA στην αγορά υπολογιστών σήμερα δεν είναι τόσο ισχυρές, γι' αυτό η δοκιμή μας εξακολουθεί να ισχύει .

Για μια εικόνα 2048 x 2048 pixel, έχουμε τα ακόλουθα αποτελέσματα.

• CPU 1 νήμα: 1.419 ms;
• CPU 2 νήματα: 749 ms;
• CPU 4 νήματα: 593 ms
• GPU (8600M GT) μπλοκ 256 νημάτων: 109 ms;
• GPU (8600M GT) μπλοκ 128 νημάτων: 94 ms;
• GPU (8800 GTX) μπλοκ 128 νημάτων/256 νημάτων: 31 ms.

Από τα αποτελέσματα μπορούν να εξαχθούν αρκετά συμπεράσματα. Ας ξεκινήσουμε με το γεγονός ότι, παρά τη συζήτηση για την εμφανή τεμπελιά των προγραμματιστών, τροποποιήσαμε την αρχική έκδοση της CPU για πολλαπλά νήματα. Όπως αναφέραμε ήδη, ο κώδικας είναι ιδανικός για αυτήν την κατάσταση - το μόνο που απαιτείται είναι να χωρίσετε την αρχική εικόνα σε όσες ζώνες υπάρχουν νήματα. Λάβετε υπόψη ότι η μετάβαση από ένα νήμα σε δύο στη διπύρηνη CPU μας είχε ως αποτέλεσμα σχεδόν γραμμική επιτάχυνση, η οποία υποδηλώνει επίσης την παράλληλη φύση του προγράμματος δοκιμής. Εντελώς απροσδόκητα, η έκδοση με τέσσερα νήματα αποδείχθηκε επίσης ταχύτερη, αν και στον επεξεργαστή μας αυτό είναι πολύ περίεργο - αντίθετα, θα μπορούσαμε να περιμένουμε πτώση της απόδοσης λόγω της επιβάρυνσης της διαχείρισης πρόσθετων νημάτων. Πώς μπορεί να εξηγηθεί αυτό το αποτέλεσμα; Είναι δύσκολο να πούμε, αλλά μπορεί να φταίει ο προγραμματιστής νημάτων των Windows. σε κάθε περίπτωση επαναλαμβάνουμε το αποτέλεσμα. Με μικρότερες υφές (512x512), το κέρδος από το νήμα δεν ήταν τόσο έντονο (περίπου 35% έναντι 100%) και η συμπεριφορά της έκδοσης με τέσσερα νήματα ήταν πιο λογική, χωρίς αύξηση σε σύγκριση με την έκδοση με δύο νήματα. Η GPU ήταν ακόμα πιο γρήγορη, αλλά όχι τόσο αισθητά ταχύτερη (το 8600M GT ήταν τρεις φορές πιο γρήγορο από την έκδοση dual-thread).

Κάντε κλικ στην εικόνα για μεγέθυνση.

Η δεύτερη σημαντική παρατήρηση είναι ότι ακόμη και η πιο αργή υλοποίηση GPU ήταν σχεδόν έξι φορές ταχύτερη από την έκδοση CPU με την υψηλότερη απόδοση. Για το πρώτο πρόγραμμα και την μη βελτιστοποιημένη έκδοση του αλγορίθμου, το αποτέλεσμα είναι πολύ ενθαρρυντικό. Λάβετε υπόψη ότι επιτύχαμε αισθητά καλύτερα αποτελέσματα σε μικρά μπλοκ, αν και η διαίσθηση μπορεί να προτείνει διαφορετικά. Η εξήγηση είναι απλή - το πρόγραμμά μας χρησιμοποιεί 14 καταχωρητές ανά νήμα και με μπλοκ 256 νημάτων, απαιτούνται 3.584 καταχωρητές ανά μπλοκ και απαιτούνται 768 νήματα για την πλήρη φόρτωση του επεξεργαστή, όπως δείξαμε. Στην περίπτωσή μας, αυτό ανέρχεται σε τρία μπλοκ ή 10.572 καταχωρητές. Αλλά ο πολυεπεξεργαστής έχει μόνο 8.192 καταχωρητές, επομένως μπορεί να κρατήσει ενεργά μόνο δύο μπλοκ. Αντίθετα, με μπλοκ 128 νημάτων, χρειαζόμαστε 1.792 καταχωρητές ανά μπλοκ. Εάν το 8.192 διαιρεθεί με το 1.792 και στρογγυλοποιηθεί στον πλησιέστερο ακέραιο, παίρνουμε τέσσερα μπλοκ. Στην πράξη, ο αριθμός των νημάτων θα είναι ο ίδιος (512 ανά πολυεπεξεργαστή, αν και θεωρητικά χρειάζονται 768 για πλήρες φορτίο), αλλά η αύξηση του αριθμού των μπλοκ δίνει στην GPU το πλεονέκτημα της ευελιξίας στην πρόσβαση στη μνήμη - όταν υπάρχει λειτουργία με μεγάλες καθυστερήσεις, μπορείτε να ξεκινήσετε την εκτέλεση εντολών από άλλο μπλοκ ενώ περιμένετε την παραλαβή των αποτελεσμάτων. Τέσσερα μπλοκ μειώνουν σαφώς την καθυστέρηση, ειδικά επειδή το πρόγραμμά μας χρησιμοποιεί πολλαπλές προσβάσεις στη μνήμη.

Ανάλυση

Τέλος, παρά τα όσα είπαμε παραπάνω, δεν μπορέσαμε να αντισταθούμε στον πειρασμό και τρέξαμε το πρόγραμμα στην 8800 GTX, η οποία ήταν τρεις φορές πιο γρήγορη από την 8600, ανεξαρτήτως μεγέθους μπλοκ. Ίσως πιστεύετε ότι στην πράξη στις αντίστοιχες αρχιτεκτονικές το αποτέλεσμα θα ήταν τέσσερις ή περισσότερες φορές υψηλότερο: 128 επεξεργαστές ALU/shader έναντι 32 και υψηλότερη ταχύτητα ρολογιού (1,35 GHz έναντι 950 MHz), αλλά δεν λειτούργησε έτσι. Πιθανότατα, η πρόσβαση στη μνήμη ήταν ο περιοριστικός παράγοντας. Για να είμαστε πιο ακριβείς, η αρχική εικόνα προσεγγίζεται ως ένας πολυδιάστατος πίνακας CUDA - ένας μάλλον περίπλοκος όρος για κάτι που δεν είναι τίποτα άλλο από μια υφή. Υπάρχουν όμως αρκετά πλεονεκτήματα.

• Οι προσβάσεις επωφελούνται από την κρυφή μνήμη υφής.
• χρησιμοποιούμε τη λειτουργία αναδίπλωσης, στην οποία δεν υπάρχει ανάγκη επεξεργασίας ορίων εικόνας, σε αντίθεση με την έκδοση της CPU.

Επιπλέον, μπορούμε να επωφεληθούμε από το "δωρεάν" φιλτράρισμα με κανονικοποιημένη διευθυνσιοδότηση μεταξύ και αντ' αυτού, αλλά στην περίπτωσή μας αυτό είναι απίθανο να είναι χρήσιμο. Όπως γνωρίζετε, το 8600 έχει 16 μονάδες υφής, σε σύγκριση με 32 για το 8800 GTX. Επομένως, η αναλογία μεταξύ των δύο αρχιτεκτονικών είναι μόνο δύο προς ένα. Προσθέστε σε αυτό τη διαφορά στις συχνότητες και παίρνουμε μια αναλογία (32 x 0,575) / (16 x 0,475) = 2,4 - κοντά στην αναλογία "τρία προς ένα" που είχαμε στην πραγματικότητα. Αυτή η θεωρία εξηγεί επίσης γιατί το μέγεθος του μπλοκ δεν αλλάζει πολύ στο G80, αφού το ALU εξακολουθεί να στηρίζεται στα μπλοκ υφής.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Εκτός από τα πολλά υποσχόμενα αποτελέσματα, η πρώτη μας γνωριμία με το CUDA πήγε πολύ καλά, δεδομένων των όχι και των πιο ευνοϊκών συνθηκών που επιλέχθηκαν. Η ανάπτυξη σε φορητό υπολογιστή Vista σημαίνει ότι θα πρέπει να χρησιμοποιήσετε το CUDA SDK 2.0, ακόμα σε έκδοση beta, με πρόγραμμα οδήγησης 174.55, το οποίο είναι επίσης beta. Παρόλα αυτά, δεν μπορούμε να αναφέρουμε δυσάρεστες εκπλήξεις - μόνο αρχικά σφάλματα κατά τον πρώτο εντοπισμό σφαλμάτων, όταν το πρόγραμμα που εξακολουθεί να παρουσιάζει σφάλματα προσπάθησε να αντιμετωπίσει τη μνήμη έξω από τον εκχωρημένο χώρο.

Η οθόνη άρχισε να τρεμοπαίζει άγρια, μετά η οθόνη μαύρισε... μέχρι που τα Vista έτρεξαν το service επισκευής του προγράμματος οδήγησης και όλα ήταν καλά. Αλλά εξακολουθεί να είναι κάπως περίεργο να το βλέπετε αυτό εάν έχετε συνηθίσει να βλέπετε ένα τυπικό σφάλμα τμηματοποίησης σε τυπικά προγράμματα όπως το δικό μας. Τέλος, μια μικρή κριτική προς την nVidia: σε όλη την τεκμηρίωση που είναι διαθέσιμη για το CUDA, δεν υπάρχει κανένας μικρός οδηγός που θα σας έλεγε βήμα προς βήμα πώς να δημιουργήσετε ένα περιβάλλον ανάπτυξης για το Visual Studio. Στην πραγματικότητα, αυτό δεν είναι μεγάλο πρόβλημα, καθώς το SDK έχει ένα πλήρες σύνολο παραδειγμάτων που μπορείτε να μελετήσετε για να κατανοήσετε το πλαίσιο για τις εφαρμογές CUDA, αλλά ένας οδηγός για αρχάριους θα ήταν χρήσιμος.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Η Nvidia παρουσίασε το CUDA με την κυκλοφορία του GeForce 8800. Και εκείνη την εποχή οι υποσχέσεις φαίνονταν πολύ δελεαστικές, αλλά συγκρατήσαμε τον ενθουσιασμό μας μέχρι να το δοκιμάσουμε πραγματικά. Πράγματι, εκείνη την εποχή φαινόταν περισσότερο σαν να μαρκάρεις την περιοχή για να παραμείνει στο κύμα GPGPU. Χωρίς ένα διαθέσιμο SDK, είναι δύσκολο να πούμε ότι αυτό δεν είναι απλώς άλλο ένα μάρκετινγκ που δεν θα λειτουργήσει. Δεν είναι η πρώτη φορά που μια καλή πρωτοβουλία ανακοινώνεται πολύ νωρίς και δεν έχει δει το φως της δημοσιότητας εκείνη την εποχή λόγω έλλειψης υποστήριξης - ειδικά σε έναν τόσο ανταγωνιστικό τομέα. Τώρα, ενάμιση χρόνο μετά την ανακοίνωση, μπορούμε με βεβαιότητα να πούμε ότι η nVidia κράτησε τον λόγο της.

Το SDK εμφανίστηκε γρήγορα σε beta στις αρχές του 2007 και έκτοτε ενημερώθηκε γρήγορα, γεγονός που αποδεικνύει τη σημασία αυτού του έργου για την nVidia. Σήμερα, το CUDA αναπτύσσεται πολύ όμορφα: το SDK είναι ήδη διαθέσιμο σε beta έκδοση 2.0 για μεγάλα λειτουργικά συστήματα (Windows XP και Vista, Linux, καθώς και 1.1 για Mac OS X) και η nVidia έχει αφιερώσει μια ολόκληρη ενότητα του ιστότοπου για προγραμματιστές.

Σε πιο επαγγελματικό επίπεδο, η εντύπωση από τα πρώτα βήματα με το CUDA αποδείχθηκε πολύ θετική. Ακόμα κι αν είστε εξοικειωμένοι με την αρχιτεκτονική GPU, μπορείτε εύκολα να το καταλάβετε. Όταν ένα API φαίνεται ξεκάθαρο με την πρώτη ματιά, αρχίζεις αμέσως να πιστεύεις ότι θα έχεις πειστικά αποτελέσματα. Δεν θα χαθεί όμως ο υπολογιστικός χρόνος από τις πολυάριθμες μεταφορές από την CPU στην GPU; Και πώς να χρησιμοποιήσετε αυτά τα χιλιάδες νήματα χωρίς ουσιαστικά κανένα πρωτόγονο συγχρονισμό; Ξεκινήσαμε τα πειράματά μας έχοντας κατά νου όλες αυτές τις ανησυχίες. Αλλά γρήγορα εξαφανίστηκαν όταν η πρώτη έκδοση του αλγορίθμου μας, αν και πολύ ασήμαντη, αποδείχθηκε ότι ήταν σημαντικά ταχύτερη από ό,τι στη CPU.

Έτσι, το CUDA δεν είναι ένα μαγικό ραβδί για ερευνητές που θέλουν να πείσουν τη διοίκηση του πανεπιστημίου να τους αγοράσει μια GeForce. Το CUDA είναι ήδη μια πλήρως προσβάσιμη τεχνολογία που μπορεί να χρησιμοποιηθεί από οποιονδήποτε προγραμματιστή με γνώση C, εάν είναι πρόθυμος να αφιερώσει χρόνο και προσπάθεια για να συνηθίσει στο νέο πρότυπο προγραμματισμού. Αυτή η προσπάθεια δεν θα πάει χαμένη εάν οι αλγόριθμοί σας είναι καλά παραλληλισμένοι. Θα θέλαμε επίσης να ευχαριστήσουμε τη nVidia για την παροχή ολοκληρωμένης και υψηλής ποιότητας τεκμηρίωσης όπου οι αρχάριοι προγραμματιστές CUDA θα βρουν απαντήσεις.

Τι χρειάζεται το CUDA για να γίνει ένα αναγνωρίσιμο API; Με μια λέξη: φορητότητα. Γνωρίζουμε ότι το μέλλον της πληροφορικής βρίσκεται στους παράλληλους υπολογιστές - σήμερα όλοι προετοιμάζονται ήδη για τέτοιες αλλαγές και όλες οι πρωτοβουλίες, τόσο λογισμικού όσο και υλικού, στοχεύουν προς αυτή την κατεύθυνση. Ωστόσο, αυτή τη στιγμή, αν κοιτάξουμε την ανάπτυξη παραδειγμάτων, βρισκόμαστε ακόμη στο αρχικό στάδιο: δημιουργούμε νήματα με μη αυτόματο τρόπο και προσπαθούμε να σχεδιάσουμε την πρόσβαση σε κοινόχρηστους πόρους. Όλα αυτά μπορούν με κάποιο τρόπο να αντιμετωπιστούν εάν ο αριθμός των πυρήνων μπορεί να μετρηθεί στα δάχτυλα του ενός χεριού. Αλλά σε λίγα χρόνια, όταν ο αριθμός των επεξεργαστών θα είναι εκατοντάδες, αυτή η δυνατότητα δεν θα υπάρχει πλέον. Με την κυκλοφορία του CUDA, η nVidia έκανε το πρώτο βήμα για την επίλυση αυτού του προβλήματος - αλλά, φυσικά, αυτή η λύση είναι κατάλληλη μόνο για GPU από αυτήν την εταιρεία και ακόμη και τότε όχι για όλους. Μόνο τα GF8 και 9 (και τα παράγωγά τους Quadro/Tesla) μπορούν να εκτελούν προγράμματα CUDA σήμερα. Και η νέα σειρά 260/280, φυσικά.

Κάντε κλικ στην εικόνα για μεγέθυνση.

Η Nvidia μπορεί να υπερηφανεύεται ότι έχει πουλήσει 70 εκατομμύρια GPU συμβατές με CUDA σε όλο τον κόσμο, αλλά αυτό εξακολουθεί να μην είναι αρκετό για να γίνει το de facto πρότυπο. Λαμβάνοντας υπόψη το γεγονός ότι οι ανταγωνιστές δεν κάθονται με σταυρωμένα τα χέρια. Η AMD προσφέρει το δικό της SDK (Stream Computing) και η Intel ανακοίνωσε μια λύση (Ct), αν και δεν είναι ακόμη διαθέσιμη. Έρχεται πόλεμος προτύπων και σαφώς δεν θα υπάρχει χώρος στην αγορά για τρεις ανταγωνιστές έως ότου ένας άλλος παίκτης, όπως η Microsoft, βγει με ένα κοινό API, το οποίο, φυσικά, θα κάνει τη ζωή πιο εύκολη για τους προγραμματιστές.

Ως εκ τούτου, η nVidia αντιμετωπίζει πολλές δυσκολίες στην έγκριση του CUDA. Αν και τεχνολογικά βρισκόμαστε αναμφίβολα αντιμέτωποι με μια επιτυχημένη λύση, εξακολουθεί να μένει να πείσουμε τους προγραμματιστές για τις προοπτικές της - και αυτό δεν θα είναι εύκολο. Ωστόσο, αν κρίνουμε από πολλές πρόσφατες ανακοινώσεις και νέα σχετικά με το API, το μέλλον δεν φαίνεται ζοφερό.

Οι συσκευές για τη μετατροπή των προσωπικών υπολογιστών σε μικρούς υπερυπολογιστές είναι γνωστές εδώ και αρκετό καιρό. Στη δεκαετία του '80 του περασμένου αιώνα, προσφέρθηκαν στην αγορά οι λεγόμενοι transputers, οι οποίοι εισήχθησαν στις τότε κοινές υποδοχές επέκτασης ISA. Στην αρχή, η απόδοσή τους σε σχετικές εργασίες ήταν εντυπωσιακή, αλλά στη συνέχεια η αύξηση της ταχύτητας των καθολικών επεξεργαστών επιταχύνθηκε, ενίσχυσαν τη θέση τους στον παράλληλο υπολογισμό και δεν υπήρχε λόγος στους πομποδέκτες. Αν και παρόμοιες συσκευές εξακολουθούν να υπάρχουν σήμερα - πρόκειται για μια ποικιλία εξειδικευμένων επιταχυντών. Συχνά όμως το πεδίο εφαρμογής τους είναι στενό και τέτοιοι επιταχυντές δεν είναι ιδιαίτερα διαδεδομένοι.

Πρόσφατα όμως, η σκυτάλη των παράλληλων υπολογιστών πέρασε στη μαζική αγορά, με τον ένα ή τον άλλο τρόπο συνδεδεμένη με τα τρισδιάστατα παιχνίδια. Οι συσκευές γενικής χρήσης με επεξεργαστές πολλαπλών πυρήνων για παράλληλους διανυσματικούς υπολογισμούς που χρησιμοποιούνται σε γραφικά 3D επιτυγχάνουν υψηλή απόδοση που δεν μπορούν να ταιριάξουν οι επεξεργαστές γενικής χρήσης. Φυσικά, η μέγιστη ταχύτητα επιτυγχάνεται μόνο σε μια σειρά από βολικές εργασίες και έχει ορισμένους περιορισμούς, αλλά τέτοιες συσκευές έχουν ήδη αρχίσει να χρησιμοποιούνται ευρέως σε περιοχές για τις οποίες δεν προορίζονταν αρχικά. Ένα εξαιρετικό παράδειγμα ενός τέτοιου παράλληλου επεξεργαστή είναι ο επεξεργαστής Cell, που αναπτύχθηκε από τη συμμαχία Sony-Toshiba-IBM και χρησιμοποιείται στην κονσόλα παιχνιδιών Sony PlayStation 3, καθώς και όλες οι σύγχρονες κάρτες βίντεο από τους ηγέτες της αγοράς - Nvidia και AMD.

Δεν θα αγγίξουμε το Cell σήμερα, αν και εμφανίστηκε νωρίτερα και είναι ένας γενικός επεξεργαστής με πρόσθετες διανυσματικές δυνατότητες, δεν το συζητάμε σήμερα. Για τους επιταχυντές τρισδιάστατων βίντεο, πριν από λίγα χρόνια εμφανίστηκαν οι πρώτες τεχνολογίες για μη γραφικούς υπολογισμούς γενικής χρήσης GPGPU (υπολογισμός γενικού σκοπού σε GPU). Εξάλλου, τα σύγχρονα τσιπ βίντεο περιέχουν εκατοντάδες μαθηματικές μονάδες εκτέλεσης και αυτή η ισχύς μπορεί να χρησιμοποιηθεί για να επιταχύνει σημαντικά πολλές υπολογιστικά εντατικές εφαρμογές. Και οι τρέχουσες γενιές των GPU έχουν μια αρκετά ευέλικτη αρχιτεκτονική, η οποία, μαζί με γλώσσες προγραμματισμού υψηλού επιπέδου και αρχιτεκτονικές υλικού-λογισμικού, όπως αυτή που συζητείται σε αυτό το άρθρο, αποκαλύπτει αυτές τις δυνατότητες και τις κάνει πολύ πιο προσιτές.

Η δημιουργία του GPCPU προκλήθηκε από την εμφάνιση αρκετά γρήγορων και ευέλικτων προγραμμάτων shader που μπορούν να εκτελεστούν από σύγχρονα τσιπ βίντεο. Οι προγραμματιστές αποφάσισαν να κάνουν τις GPU να υπολογίζουν όχι μόνο εικόνες σε τρισδιάστατες εφαρμογές, αλλά να χρησιμοποιούνται και σε άλλους παράλληλους υπολογισμούς. Στο GPGPU, χρησιμοποιήθηκαν API γραφικών για αυτό: OpenGL και Direct3D, όταν τα δεδομένα μεταφέρονταν στο τσιπ βίντεο με τη μορφή υφών και τα προγράμματα υπολογισμού φορτώθηκαν με τη μορφή shaders. Τα μειονεκτήματα αυτής της μεθόδου είναι η σχετικά υψηλή πολυπλοκότητα του προγραμματισμού, η χαμηλή ταχύτητα ανταλλαγής δεδομένων μεταξύ της CPU και της GPU και άλλοι περιορισμοί, τους οποίους θα συζητήσουμε αργότερα.

Ο υπολογισμός GPU έχει αναπτυχθεί και αναπτύσσεται πολύ γρήγορα. Και στη συνέχεια, δύο μεγάλοι κατασκευαστές τσιπ βίντεο, η Nvidia και η AMD, ανέπτυξαν και ανακοίνωσαν αντίστοιχες πλατφόρμες που ονομάζονται CUDA (Compute Unified Device Architecture) και CTM (Close To Metal ή AMD Stream Computing), αντίστοιχα. Σε αντίθεση με τα προηγούμενα μοντέλα προγραμματισμού GPU, αυτά σχεδιάστηκαν με άμεση πρόσβαση στις δυνατότητες υλικού των καρτών βίντεο. Οι πλατφόρμες δεν είναι συμβατές μεταξύ τους, η CUDA είναι μια επέκταση της γλώσσας προγραμματισμού C και η CTM είναι μια εικονική μηχανή που εκτελεί κώδικα συναρμολόγησης. Αλλά και οι δύο πλατφόρμες εξάλειψαν ορισμένους από τους σημαντικούς περιορισμούς των προηγούμενων μοντέλων GPGPU που χρησιμοποιούσαν μια παραδοσιακή διοχέτευση γραφικών και τις αντίστοιχες διεπαφές Direct3D ή OpenGL.

Φυσικά, τα ανοιχτά πρότυπα που χρησιμοποιούν OpenGL φαίνεται να είναι τα πιο φορητά και καθολικά, επιτρέποντας τη χρήση του ίδιου κώδικα σε τσιπ βίντεο διαφορετικών κατασκευαστών. Αλλά τέτοιες μέθοδοι έχουν πολλά μειονεκτήματα, είναι πολύ λιγότερο ευέλικτες και δεν είναι τόσο βολικές στη χρήση. Επιπλέον, δεν επιτρέπουν τη χρήση συγκεκριμένων δυνατοτήτων ορισμένων καρτών γραφικών, όπως η γρήγορη κοινή χρήση (κοινή) μνήμη που υπάρχει στους σύγχρονους επεξεργαστές υπολογιστών.

Αυτός είναι ο λόγος για τον οποίο η Nvidia κυκλοφόρησε την πλατφόρμα CUDA, μια γλώσσα προγραμματισμού τύπου C με δικό της μεταγλωττιστή και βιβλιοθήκες για υπολογιστές GPU. Φυσικά, η σύνταξη βέλτιστου κώδικα για τσιπ βίντεο δεν είναι καθόλου τόσο απλή και αυτή η εργασία απαιτεί πολλή χειρωνακτική εργασία, αλλά το CUDA αποκαλύπτει όλες τις δυνατότητες και δίνει στον προγραμματιστή περισσότερο έλεγχο στις δυνατότητες υλικού της GPU. Είναι σημαντικό ότι η υποστήριξη Nvidia CUDA είναι διαθέσιμη στα τσιπ G8x, G9x και GT2xx που χρησιμοποιούνται στις κάρτες γραφικών της σειράς Geforce 8, 9 και 200, οι οποίες είναι πολύ διαδεδομένες. Τώρα κυκλοφόρησε η τελική έκδοση του CUDA 2.0, η οποία φέρνει ορισμένα νέα χαρακτηριστικά, όπως υποστήριξη για υπολογισμούς διπλής ακρίβειας. Το CUDA είναι διαθέσιμο σε λειτουργικά συστήματα Linux, 32-bit και 64-bit Linux, Windows και MacOS X.

Διαφορά μεταξύ CPU και GPU σε παράλληλους υπολογισμούς

Η αύξηση των συχνοτήτων των καθολικών επεξεργαστών έχει έρθει αντιμέτωπη με φυσικούς περιορισμούς και υψηλή κατανάλωση ενέργειας και η απόδοσή τους αυξάνεται ολοένα και περισσότερο με την τοποθέτηση πολλών πυρήνων σε ένα τσιπ. Οι επεξεργαστές που πωλούνται αυτήν τη στιγμή περιέχουν μόνο έως τέσσερις πυρήνες (η περαιτέρω ανάπτυξη δεν θα είναι γρήγορη) και είναι σχεδιασμένοι για γενικές εφαρμογές, χρησιμοποιώντας πολλαπλές εντολές και ροή δεδομένων MIMD. Κάθε πυρήνας λειτουργεί ξεχωριστά από τους άλλους, εκτελώντας διαφορετικές εντολές για διαφορετικές διεργασίες.

Εξειδικευμένες διανυσματικές δυνατότητες (SSE2 και SSE3) για διανύσματα τεσσάρων συστατικών (μονής ακρίβειας κινητής υποδιαστολής) και δύο συστατικών (διπλής ακρίβειας) εμφανίστηκαν σε επεξεργαστές γενικής χρήσης λόγω των αυξημένων απαιτήσεων των εφαρμογών γραφικών, κυρίως. Αυτός είναι ο λόγος για τον οποίο για ορισμένες εργασίες η χρήση των GPU είναι πιο κερδοφόρα, επειδή αρχικά είχαν κατασκευαστεί για αυτές.

Για παράδειγμα, στα τσιπ βίντεο Nvidia η κύρια μονάδα είναι ένας πολυεπεξεργαστής με οκτώ έως δέκα πυρήνες και εκατοντάδες ALU συνολικά, αρκετές χιλιάδες καταχωρητές και μια μικρή ποσότητα κοινής μνήμης. Επιπλέον, η κάρτα βίντεο περιέχει γρήγορη παγκόσμια μνήμη με πρόσβαση σε αυτήν από όλους τους πολυεπεξεργαστές, τοπική μνήμη σε κάθε πολυεπεξεργαστή και ειδική μνήμη για σταθερές.

Το πιο σημαντικό, αυτοί οι πολλαπλοί πυρήνες πολλαπλών επεξεργαστών στη GPU είναι πυρήνες SIMD (μονή ροή εντολών, πολλαπλή ροή δεδομένων). Και αυτοί οι πυρήνες εκτελούν τις ίδιες οδηγίες ταυτόχρονα, ένα στυλ προγραμματισμού που είναι κοινό για αλγόριθμους γραφικών και πολλές επιστημονικές εργασίες, αλλά απαιτεί συγκεκριμένο προγραμματισμό. Αλλά αυτή η προσέγγιση σάς επιτρέπει να αυξήσετε τον αριθμό των μονάδων εκτέλεσης λόγω της απλοποίησής τους.

Λοιπόν, ας απαριθμήσουμε τις κύριες διαφορές μεταξύ των αρχιτεκτονικών CPU και GPU. Οι πυρήνες της CPU έχουν σχεδιαστεί για να εκτελούν μία μόνο ροή διαδοχικών εντολών με μέγιστη απόδοση, ενώ οι GPU έχουν σχεδιαστεί για να εκτελούν γρήγορα μεγάλο αριθμό παράλληλων ροών εντολών. Οι επεξεργαστές γενικής χρήσης είναι βελτιστοποιημένοι για να επιτυγχάνουν υψηλή απόδοση από μία μόνο ροή εντολών, επεξεργάζοντας τόσο ακέραιους όσο και αριθμούς κινητής υποδιαστολής. Σε αυτήν την περίπτωση, η πρόσβαση στη μνήμη είναι τυχαία.

Οι προγραμματιστές CPU προσπαθούν να εκτελέσουν όσο το δυνατόν περισσότερες εντολές παράλληλα για να αυξήσουν την απόδοση. Για να επιτευχθεί αυτό, ξεκινώντας με τους επεξεργαστές Intel Pentium, εμφανίστηκε η υπερκλιμακωτή εκτέλεση, εξασφαλίζοντας την εκτέλεση δύο εντολών ανά κύκλο ρολογιού και το Pentium Pro διακρίθηκε από την ασυνήθιστη εκτέλεση εντολών. Όμως η παράλληλη εκτέλεση μιας διαδοχικής ροής εντολών έχει ορισμένους βασικούς περιορισμούς και η αύξηση του αριθμού των μονάδων εκτέλεσης δεν μπορεί να επιτύχει πολλαπλή αύξηση της ταχύτητας.

Τα τσιπ βίντεο έχουν μια απλή και παράλληλη λειτουργία από την αρχή. Το τσιπ βίντεο παίρνει μια ομάδα πολυγώνων ως είσοδο, εκτελεί όλες τις απαραίτητες λειτουργίες και παράγει pixel ως έξοδο. Η επεξεργασία των πολυγώνων και των εικονοστοιχείων είναι ανεξάρτητη. Επομένως, λόγω της εγγενώς παράλληλης οργάνωσης της εργασίας, η GPU χρησιμοποιεί μεγάλο αριθμό μονάδων εκτέλεσης, οι οποίες είναι εύκολο να φορτωθούν, σε αντίθεση με τη διαδοχική ροή εντολών για την CPU. Επιπλέον, οι σύγχρονες GPU μπορούν επίσης να εκτελούν περισσότερες από μία εντολές ανά κύκλο ρολογιού (διπλή έκδοση). Έτσι, η αρχιτεκτονική Tesla, υπό ορισμένες προϋποθέσεις, εκκινεί ταυτόχρονα λειτουργίες MAD+MUL ή MAD+SFU.

Η GPU διαφέρει επίσης από την CPU όσον αφορά τις αρχές πρόσβασης στη μνήμη. Στη GPU είναι συνεκτικό και εύκολα προβλέψιμο - εάν ένα texture texel διαβάζεται από τη μνήμη, τότε μετά από λίγο θα έρθει η ώρα για γειτονικά texel. Και κατά την εγγραφή, συμβαίνει το ίδιο - ένα pixel γράφεται στο framebuffer και μετά από μερικούς κύκλους ρολογιού θα εγγραφεί αυτό που βρίσκεται δίπλα του. Επομένως, η οργάνωση της μνήμης είναι διαφορετική από αυτή που χρησιμοποιείται στην CPU. Και το τσιπ βίντεο, σε αντίθεση με τους καθολικούς επεξεργαστές, απλά δεν χρειάζεται μεγάλη μνήμη cache και οι υφές απαιτούν μόνο λίγα (έως 128-256 στις τρέχουσες GPU) kilobyte.

Και η ίδια η εργασία με τη μνήμη είναι κάπως διαφορετική για τις GPU και τις CPU. Έτσι, δεν έχουν όλοι οι κεντρικοί επεξεργαστές ενσωματωμένους ελεγκτές μνήμης και όλες οι GPU έχουν συνήθως πολλούς ελεγκτές, έως και οκτώ κανάλια 64-bit στο τσιπ Nvidia GT200. Επιπλέον, οι κάρτες γραφικών χρησιμοποιούν ταχύτερη μνήμη και ως εκ τούτου, τα τσιπ βίντεο έχουν πρόσβαση σε σημαντικά μεγαλύτερο εύρος ζώνης μνήμης, το οποίο είναι επίσης πολύ σημαντικό για παράλληλους υπολογισμούς που λειτουργούν με τεράστιες ροές δεδομένων.

Στους γενικούς επεξεργαστές, μεγάλος αριθμός τρανζίστορ και περιοχή τσιπ χρησιμοποιούνται για buffer εντολών, πρόβλεψη κλάδου υλικού και τεράστιες ποσότητες μνήμης cache στο τσιπ. Όλα αυτά τα μπλοκ υλικού χρειάζονται για να επιταχυνθεί η εκτέλεση μερικών ροών εντολών. Τα τσιπ βίντεο ξοδεύουν τρανζίστορ σε συστοιχίες μονάδων εκτέλεσης, μονάδες ελέγχου ροής, μικρή κοινόχρηστη μνήμη και ελεγκτές μνήμης για πολλά κανάλια. Τα παραπάνω δεν επιταχύνουν την εκτέλεση μεμονωμένων νημάτων, αλλά επιτρέπουν στο τσιπ να χειρίζεται πολλές χιλιάδες νήματα που εκτελούνται ταυτόχρονα από το τσιπ και απαιτούν υψηλό εύρος ζώνης μνήμης.

Σχετικά με τις διαφορές στην προσωρινή αποθήκευση. Οι CPU γενικής χρήσης χρησιμοποιούν κρυφή μνήμη για να αυξήσουν την απόδοση μειώνοντας την καθυστέρηση πρόσβασης στη μνήμη, ενώ οι GPU χρησιμοποιούν κρυφή μνήμη ή κοινόχρηστη μνήμη για να αυξήσουν το εύρος ζώνης. Οι CPU μειώνουν την καθυστέρηση πρόσβασης στη μνήμη χρησιμοποιώντας μεγάλες κρυφές μνήμες και πρόβλεψη διακλάδωσης κώδικα. Αυτά τα κομμάτια υλικού καταλαμβάνουν το μεγαλύτερο μέρος της περιοχής του τσιπ και καταναλώνουν πολλή ισχύ. Τα τσιπ βίντεο παρακάμπτουν το πρόβλημα των καθυστερήσεων πρόσβασης στη μνήμη εκτελώντας χιλιάδες νήματα ταυτόχρονα - ενώ ένα από τα νήματα περιμένει δεδομένα από τη μνήμη, το τσιπ βίντεο μπορεί να εκτελέσει υπολογισμούς από άλλο νήμα χωρίς αναμονή ή καθυστερήσεις.

Υπάρχουν επίσης πολλές διαφορές στην υποστήριξη πολλαπλών νημάτων. Η CPU εκτελεί 1-2 νήματα υπολογισμών ανά πυρήνα επεξεργαστή και τα τσιπ βίντεο μπορούν να υποστηρίξουν έως και 1024 νήματα ανά πολυεπεξεργαστή, από τα οποία υπάρχουν πολλά στο τσιπ. Και αν η εναλλαγή από το ένα νήμα στο άλλο κοστίζει εκατοντάδες κύκλους ρολογιού για την CPU, τότε η GPU αλλάζει πολλά νήματα σε έναν κύκλο ρολογιού.

Επιπλέον, οι CPU χρησιμοποιούν μονάδες SIMD (μονής εντολής σε πολλαπλά δεδομένα) για διανυσματικούς υπολογισμούς και οι GPU χρησιμοποιούν SIMT (πολλά νήματα μιας εντολής) για επεξεργασία βαθμωτών νημάτων. Το SIMT δεν απαιτεί από τον προγραμματιστή να μετατρέπει δεδομένα σε διανύσματα και επιτρέπει αυθαίρετη διακλάδωση σε ροές.

Εν ολίγοις, μπορούμε να πούμε ότι, σε αντίθεση με τις σύγχρονες καθολικές CPU, τα τσιπ βίντεο έχουν σχεδιαστεί για παράλληλους υπολογισμούς με μεγάλο αριθμό αριθμητικών πράξεων. Και ένας σημαντικά μεγαλύτερος αριθμός τρανζίστορ GPU λειτουργούν για τον προορισμό τους - την επεξεργασία συστοιχιών δεδομένων και δεν ελέγχουν την εκτέλεση (έλεγχος ροής) μερικών διαδοχικών υπολογιστικών νημάτων. Αυτό είναι ένα διάγραμμα του πόσο χώρο καταλαμβάνουν διάφορες λογικές στη CPU και την GPU:

Ως αποτέλεσμα, η βάση για την αποτελεσματική χρήση της ισχύος της GPU σε επιστημονικούς και άλλους μη γραφικούς υπολογισμούς είναι η παραλληλοποίηση αλγορίθμων σε εκατοντάδες μονάδες εκτέλεσης που είναι διαθέσιμες σε τσιπ βίντεο. Για παράδειγμα, πολλές εφαρμογές μοριακής μοντελοποίησης είναι κατάλληλες για υπολογισμούς σε τσιπ βίντεο, απαιτούν μεγάλη υπολογιστική ισχύ και ως εκ τούτου είναι βολικές για παράλληλους υπολογιστές. Και η χρήση πολλαπλών GPU παρέχει ακόμη μεγαλύτερη υπολογιστική ισχύ για την επίλυση τέτοιων προβλημάτων.

Η εκτέλεση υπολογισμών στη GPU δείχνει εξαιρετικά αποτελέσματα σε αλγόριθμους που χρησιμοποιούν παράλληλη επεξεργασία δεδομένων. Δηλαδή, όταν η ίδια ακολουθία μαθηματικών πράξεων εφαρμόζεται σε μεγάλο όγκο δεδομένων. Σε αυτή την περίπτωση, τα καλύτερα αποτελέσματα επιτυγχάνονται εάν ο λόγος του αριθμού των αριθμητικών εντολών προς τον αριθμό των προσβάσεων στη μνήμη είναι αρκετά μεγάλος. Αυτό δημιουργεί λιγότερες απαιτήσεις στον έλεγχο ροής και η υψηλή πυκνότητα των μαθηματικών και ο μεγάλος όγκος δεδομένων εξαλείφουν την ανάγκη για μεγάλες κρυφές μνήμες, όπως στην CPU.

Ως αποτέλεσμα όλων των διαφορών που περιγράφονται παραπάνω, η θεωρητική απόδοση των τσιπ βίντεο υπερβαίνει σημαντικά την απόδοση της CPU. Η Nvidia παρέχει το ακόλουθο γράφημα της αύξησης της απόδοσης της CPU και της GPU τα τελευταία χρόνια:

Φυσικά, αυτά τα δεδομένα δεν είναι χωρίς μερίδιο δόλου. Άλλωστε, σε μια CPU είναι πολύ πιο εύκολο στην πράξη να επιτευχθούν θεωρητικά στοιχεία, και οι αριθμοί δίνονται για απλή ακρίβεια στην περίπτωση μιας GPU και για διπλή ακρίβεια στην περίπτωση της CPU. Σε κάθε περίπτωση, η απλή ακρίβεια είναι αρκετή για ορισμένες παράλληλες εργασίες και η διαφορά ταχύτητας μεταξύ των γενικών επεξεργαστών και των επεξεργαστών γραφικών είναι πολύ μεγάλη και επομένως το παιχνίδι αξίζει το κερί.

Πρώτες προσπάθειες χρήσης υπολογισμών GPU

Προσπαθούν να χρησιμοποιήσουν τσιπ βίντεο σε παράλληλους μαθηματικούς υπολογισμούς εδώ και αρκετό καιρό. Οι πρώτες απόπειρες σε τέτοιες εφαρμογές ήταν εξαιρετικά πρωτόγονες και περιορίζονταν στη χρήση ορισμένων λειτουργιών υλικού, όπως η ραστεροποίηση και η αποθήκευση σε buffer. Αλλά τον τρέχοντα αιώνα, με την εμφάνιση των shaders, οι υπολογισμοί μήτρας άρχισαν να επιταχύνονται. Το 2003, στο SIGGRAPH, διατέθηκε μια ξεχωριστή ενότητα για υπολογιστές GPU και ονομάστηκε GPGPU (υπολογισμός γενικού σκοπού σε GPU).

Ο πιο γνωστός είναι ο μεταγλωττιστής γλώσσας προγραμματισμού BrookGPU Brook stream, σχεδιασμένος να εκτελεί μη γραφικούς υπολογισμούς στη GPU. Πριν από την εμφάνισή του, οι προγραμματιστές που χρησιμοποιούν τις δυνατότητες των τσιπ βίντεο για υπολογισμούς επέλεξαν ένα από τα δύο κοινά API: Direct3D ή OpenGL. Αυτό περιόρισε σοβαρά τη χρήση των GPU, επειδή τα τρισδιάστατα γραφικά χρησιμοποιούν shaders και textures για τα οποία δεν απαιτείται να γνωρίζουν οι ειδικοί του παράλληλου προγραμματισμού χρησιμοποιούν νήματα και πυρήνες. Ο Μπρουκ μπόρεσε να τους βοηθήσει να κάνουν το έργο τους πιο εύκολο. Αυτές οι επεκτάσεις ροής στη γλώσσα C, που αναπτύχθηκαν στο Πανεπιστήμιο του Στάνφορντ, έκρυψαν το 3D API από τους προγραμματιστές και παρουσίασαν το τσιπ βίντεο ως παράλληλο συνεπεξεργαστή. Ο μεταγλωττιστής επεξεργάστηκε το αρχείο .br με κώδικα και επεκτάσεις C++, δημιουργώντας κώδικα συνδεδεμένο με μια βιβλιοθήκη με δυνατότητα DirectX, OpenGL ή x86.

Όπως ήταν φυσικό, ο Μπρουκ είχε πολλές ελλείψεις, στις οποίες σταθήκαμε και θα συζητήσουμε αναλυτικότερα στη συνέχεια. Αλλά ακόμη και μόνο η εμφάνισή του προκάλεσε σημαντική αύξηση της προσοχής από την ίδια Nvidia και ATI στην πρωτοβουλία υπολογιστών σε GPU, καθώς η ανάπτυξη αυτών των δυνατοτήτων άλλαξε σοβαρά την αγορά στο μέλλον, ανοίγοντας έναν εντελώς νέο τομέα της - παράλληλους υπολογιστές βασίζεται σε τσιπ βίντεο.

Στη συνέχεια, ορισμένοι ερευνητές από το έργο Brook εντάχθηκαν στην ομάδα ανάπτυξης της Nvidia για να εισαγάγουν μια στρατηγική παράλληλων υπολογιστών υλικού-λογισμικού, ανοίγοντας νέο μερίδιο αγοράς. Και το κύριο πλεονέκτημα αυτής της πρωτοβουλίας της Nvidia είναι ότι οι προγραμματιστές γνωρίζουν όλες τις δυνατότητες των GPU τους μέχρι την τελευταία λεπτομέρεια και δεν χρειάζεται να χρησιμοποιήσετε το API γραφικών και μπορείτε να εργαστείτε με το υλικό απευθείας χρησιμοποιώντας το πρόγραμμα οδήγησης. Το αποτέλεσμα των προσπαθειών αυτής της ομάδας ήταν το Nvidia CUDA (Compute Unified Device Architecture) μια νέα αρχιτεκτονική λογισμικού και υλικού για παράλληλους υπολογιστές σε Nvidia GPU, το οποίο είναι το θέμα αυτού του άρθρου.

Τομείς εφαρμογής παράλληλων υπολογισμών σε GPU

Για να κατανοήσετε τα πλεονεκτήματα της μεταφοράς υπολογισμών σε τσιπ βίντεο, ακολουθούν οι μέσοι αριθμοί που ελήφθησαν από ερευνητές σε όλο τον κόσμο. Κατά μέσο όρο, κατά τη μεταφορά υπολογισμών στη GPU, πολλές εργασίες επιτυγχάνουν ταχύτητες 5-30 φορές σε σύγκριση με γρήγορους επεξεργαστές γενικής χρήσης. Οι μεγαλύτεροι αριθμοί (της τάξεως της επιτάχυνσης 100x ή και περισσότερο!) επιτυγχάνονται με κώδικα που δεν είναι πολύ κατάλληλος για υπολογισμούς με χρήση μπλοκ SSE, αλλά είναι αρκετά βολικός για GPU.

Αυτά είναι μόνο μερικά παραδείγματα επιταχύνσεων για συνθετικό κώδικα στη GPU έναντι του διανυσματικού κώδικα SSE στη CPU (σύμφωνα με τη Nvidia):

• Μικροσκόπιο φθορισμού: 12x;
• Μοριακή δυναμική (μη δεσμευμένη δύναμη υπολογ.): 8-16x;
• Ηλεκτροστατική (άμεση και πολυεπίπεδη άθροιση Coulomb): 40-120x και 7x.

Και αυτό είναι ένα σημάδι ότι η Nvidia αγαπά πολύ, δείχνοντάς το σε όλες τις παρουσιάσεις, στις οποίες θα σταθούμε λεπτομερέστερα στο δεύτερο μέρος του άρθρου, αφιερωμένο σε συγκεκριμένα παραδείγματα πρακτικών εφαρμογών του CUDA computing:

Όπως μπορείτε να δείτε, τα νούμερα είναι πολύ ελκυστικά, οι αυξήσεις 100-150 φορές είναι ιδιαίτερα εντυπωσιακές. Στο επόμενο άρθρο για το CUDA, θα δούμε μερικούς από αυτούς τους αριθμούς λεπτομερώς. Ας απαριθμήσουμε τώρα τις κύριες εφαρμογές στις οποίες χρησιμοποιείται σήμερα ο υπολογισμός GPU: ανάλυση και επεξεργασία εικόνων και σημάτων, φυσική προσομοίωση, υπολογιστικά μαθηματικά, υπολογιστική βιολογία, οικονομικοί υπολογισμοί, βάσεις δεδομένων, δυναμική αερίων και υγρών, κρυπτογραφία, προσαρμοστική ακτινοθεραπεία, αστρονομία, επεξεργασία ήχου, βιοπληροφορική, βιολογικές προσομοιώσεις, όραση υπολογιστή, εξόρυξη δεδομένων, ψηφιακός κινηματογράφος και τηλεόραση, ηλεκτρομαγνητικές προσομοιώσεις, γεωγραφικά συστήματα πληροφοριών, στρατιωτικές εφαρμογές, σχεδιασμός ορυχείων, μοριακή δυναμική, απεικόνιση μαγνητικού συντονισμού (MRI), νευρωνικά δίκτυα, ωκεανογραφική έρευνα, φυσική σωματιδίων , προσομοίωση αναδίπλωσης πρωτεϊνών, κβαντική χημεία, ανίχνευση ακτίνων, οπτικοποίηση, ραντάρ, προσομοίωση δεξαμενής, τεχνητή νοημοσύνη, δορυφορική ανάλυση δεδομένων, σεισμική εξερεύνηση, χειρουργική επέμβαση, υπερηχογράφημα, τηλεδιάσκεψη.

Λεπτομέρειες για πολλές εφαρμογές μπορείτε να βρείτε στον ιστότοπο της Nvidia στην ενότητα για. Όπως μπορείτε να δείτε, η λίστα είναι αρκετά μεγάλη, αλλά δεν είναι μόνο αυτό! Μπορεί να συνεχιστεί, και σίγουρα μπορούμε να υποθέσουμε ότι στο μέλλον θα βρεθούν άλλοι τομείς εφαρμογής παράλληλων υπολογισμών σε τσιπ βίντεο, τους οποίους δεν γνωρίζουμε ακόμη.

Χαρακτηριστικά Nvidia CUDA

Η τεχνολογία CUDA είναι μια αρχιτεκτονική υπολογιστών υλικού-λογισμικού Nvidia που βασίζεται σε μια επέκταση της γλώσσας C, η οποία καθιστά δυνατή την οργάνωση της πρόσβασης στο σύνολο εντολών ενός επιταχυντή γραφικών και τη διαχείριση της μνήμης του κατά την οργάνωση παράλληλων υπολογιστών. Το CUDA βοηθά στην υλοποίηση εκτελέσιμων αλγορίθμων σε επιταχυντές βίντεο GeForce όγδοης γενιάς και παλαιότερους (Geforce 8, Geforce 9, Geforce 200 series), καθώς και σε Quadro και Tesla.

Αν και η πολυπλοκότητα του προγραμματισμού των GPU με χρήση CUDA είναι αρκετά υψηλή, είναι μικρότερη από ό,τι με προηγούμενες λύσεις GPGPU. Τέτοια προγράμματα απαιτούν τη διαίρεση της εφαρμογής μεταξύ πολλαπλών πολυεπεξεργαστών, παρόμοια με τον προγραμματισμό MPI, αλλά χωρίς κοινή χρήση των δεδομένων που είναι αποθηκευμένα στην κοινόχρηστη μνήμη βίντεο. Και επειδή ο προγραμματισμός CUDA για κάθε πολυεπεξεργαστή είναι παρόμοιος με τον προγραμματισμό OpenMP, απαιτεί καλή κατανόηση της οργάνωσης της μνήμης. Αλλά, φυσικά, η πολυπλοκότητα της ανάπτυξης και της μεταφοράς στο CUDA εξαρτάται σε μεγάλο βαθμό από την εφαρμογή.

Το κιτ προγραμματιστή περιέχει πολλά παραδείγματα κώδικα και είναι καλά τεκμηριωμένο. Η διαδικασία εκμάθησης θα διαρκέσει περίπου δύο έως τέσσερις εβδομάδες για όσους είναι ήδη εξοικειωμένοι με το OpenMP και το MPI. Το API βασίζεται στην εκτεταμένη γλώσσα C και για τη μετάφραση κώδικα από αυτήν τη γλώσσα, το CUDA SDK περιλαμβάνει τον μεταγλωττιστή γραμμής εντολών nvcc, που βασίζεται στον μεταγλωττιστή Open64 ανοιχτού κώδικα.

Ας απαριθμήσουμε τα κύρια χαρακτηριστικά του CUDA:

• ενοποιημένη λύση λογισμικού και υλικού για παράλληλους υπολογιστές σε τσιπ βίντεο Nvidia.
• μια ευρεία γκάμα υποστηριζόμενων λύσεων, από φορητές συσκευές έως multi-chip
• τυπική γλώσσα προγραμματισμού C?
• τυπικές βιβλιοθήκες αριθμητικής ανάλυσης FFT (Fast Fourier Transform) και BLAS (Γραμμική Άλγεβρα).
• βελτιστοποιημένη ανταλλαγή δεδομένων μεταξύ CPU και GPU.
• αλληλεπίδραση με τα API γραφικών OpenGL και DirectX.
• υποστήριξη για λειτουργικά συστήματα 32 και 64 bit: Windows XP, Windows Vista, Linux και MacOS X.
• Δυνατότητα ανάπτυξης σε χαμηλό επίπεδο.

Όσον αφορά την υποστήριξη λειτουργικών συστημάτων, θα πρέπει να προστεθεί ότι όλες οι μεγάλες διανομές Linux υποστηρίζονται επίσημα (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), αλλά, αν κρίνουμε από τα δεδομένα των ενθουσιωδών, το CUDA λειτουργεί εξαιρετικά σε άλλες εκδόσεις: Fedora Core, Ubuntu, Gentoo, κ.λπ.

Το περιβάλλον ανάπτυξης CUDA (CUDA Toolkit) περιλαμβάνει:

• nvcc μεταγλωττιστής?
• Βιβλιοθήκες FFT και BLAS.
• Προφίλ
• πρόγραμμα εντοπισμού σφαλμάτων gdb για GPU.
• Το πρόγραμμα οδήγησης χρόνου εκτέλεσης CUDA περιλαμβάνεται στα τυπικά προγράμματα οδήγησης Nvidia
• εγχειρίδιο προγραμματισμού?
• CUDA Developer SDK (πηγαίος κώδικας, βοηθητικά προγράμματα και τεκμηρίωση).

Παραδείγματα στον πηγαίο κώδικα: παράλληλη bitonic ταξινόμηση, μεταφορά μήτρας, άθροιση παράλληλου προθέματος μεγάλων πινάκων, συνέλιξη εικόνας, διακριτός μετασχηματισμός κυματιδίων, παράδειγμα αλληλεπίδρασης με OpenGL και Direct3D, χρήση των βιβλιοθηκών CUBLAS και CUFFT, υπολογισμός της τιμής επιλογής (Μαύρο τύπος Scholes, διωνυμικό μοντέλο, μέθοδος Monte Carlo), παράλληλη γεννήτρια τυχαίων αριθμών Mersenne Twister, υπολογισμός ιστογράμματος μεγάλου πίνακα, μείωση θορύβου, φίλτρο Sobel (εύρεση ορίων).

Πλεονεκτήματα και περιορισμοί του CUDA

Από την οπτική γωνία ενός προγραμματιστή, μια διοχέτευση γραφικών είναι μια συλλογή σταδίων επεξεργασίας. Το μπλοκ γεωμετρίας δημιουργεί τα τρίγωνα και το μπλοκ ραστεροποίησης δημιουργεί τα εικονοστοιχεία που εμφανίζονται στην οθόνη. Το παραδοσιακό μοντέλο προγραμματισμού GPGPU μοιάζει με αυτό:

Για να μεταφερθούν οι υπολογισμοί στη GPU σε αυτό το μοντέλο, απαιτείται ειδική προσέγγιση. Ακόμη και η προσθήκη δύο διανυσμάτων βάσει στοιχείων θα απαιτήσει τη σχεδίαση του σχήματος στην οθόνη ή σε μια προσωρινή μνήμη εκτός οθόνης. Το σχήμα είναι ραστεροποιημένο, το χρώμα κάθε pixel υπολογίζεται χρησιμοποιώντας ένα δεδομένο πρόγραμμα (pixel shader). Το πρόγραμμα διαβάζει τα δεδομένα εισόδου από τις υφές για κάθε pixel, τα προσθέτει και τα γράφει στο buffer εξόδου. Και όλες αυτές οι πολυάριθμες λειτουργίες χρειάζονται για κάτι που είναι γραμμένο σε έναν μόνο χειριστή σε μια κανονική γλώσσα προγραμματισμού!

Επομένως, η χρήση του GPGPU για υπολογιστές γενικού σκοπού έχει τον περιορισμό ότι είναι πολύ δύσκολη για την εκπαίδευση προγραμματιστών. Ναι, και υπάρχουν αρκετοί άλλοι περιορισμοί, επειδή ένας σκιαστής εικονοστοιχείων είναι απλώς ένας τύπος για την εξάρτηση του τελικού χρώματος ενός εικονοστοιχείου από τη συντεταγμένη του και η γλώσσα των σκιερών εικονοστοιχείων είναι μια γλώσσα για τη σύνταξη αυτών των τύπων με σύνταξη τύπου C . Οι πρώιμες μέθοδοι GPGPU είναι ένα προσεγμένο κόλπο που σας επιτρέπει να χρησιμοποιήσετε τη δύναμη της GPU, αλλά χωρίς καμία από τις ευκολίες. Τα δεδομένα εκεί αντιπροσωπεύονται από εικόνες (υφές) και ο αλγόριθμος αντιπροσωπεύεται από τη διαδικασία ραστεροποίησης. Ιδιαίτερη αναφορά είναι το πολύ συγκεκριμένο μοντέλο μνήμης και εκτέλεσης.

Η αρχιτεκτονική λογισμικού και υλικού της Nvidia για υπολογιστές GPU διαφέρει από τα προηγούμενα μοντέλα GPGPU, καθώς σας επιτρέπει να γράφετε προγράμματα για τη GPU σε πραγματική γλώσσα C με τυπική σύνταξη, δείκτες και την ανάγκη για ελάχιστες επεκτάσεις για πρόσβαση στους υπολογιστικούς πόρους των τσιπ βίντεο. Το CUDA είναι ανεξάρτητο από τα API γραφικών και έχει ορισμένες δυνατότητες που έχουν σχεδιαστεί ειδικά για υπολογιστές γενικού σκοπού.

Πλεονεκτήματα του CUDA έναντι της παραδοσιακής προσέγγισης στον υπολογισμό GPGPU:

• Η διεπαφή προγραμματισμού εφαρμογής CUDA βασίζεται στην τυπική γλώσσα προγραμματισμού C με επεκτάσεις, η οποία απλοποιεί τη διαδικασία εκμάθησης και υλοποίησης της αρχιτεκτονικής CUDA.
• Το CUDA παρέχει πρόσβαση σε 16 KB κοινόχρηστης μνήμης νημάτων ανά πολυεπεξεργαστή, η οποία μπορεί να χρησιμοποιηθεί για την οργάνωση μιας κρυφής μνήμης με μεγαλύτερο εύρος ζώνης από τις ανακτήσεις υφής.
• πιο αποτελεσματική μεταφορά δεδομένων μεταξύ συστήματος και μνήμης βίντεο
• δεν χρειάζονται γραφικά API με πλεονασμό και επιβάρυνση.
• Διευθυνσιοδότηση γραμμικής μνήμης, συγκέντρωση και διασπορά, δυνατότητα εγγραφής σε αυθαίρετες διευθύνσεις.
• υποστήριξη υλικού για λειτουργίες ακεραίων και bit.

Κύριοι περιορισμοί του CUDA:

• έλλειψη υποστήριξης αναδρομής για εκτελέσιμες συναρτήσεις.
• ελάχιστο πλάτος μπλοκ 32 νημάτων.
• κλειστή αρχιτεκτονική CUDA που ανήκει στην Nvidia.

Οι αδυναμίες του προγραμματισμού με προηγούμενες μεθόδους GPGPU είναι ότι αυτές οι μέθοδοι δεν χρησιμοποιούν μονάδες εκτέλεσης σκίασης κορυφής σε προηγούμενες μη ενοποιημένες αρχιτεκτονικές, τα δεδομένα αποθηκεύονται σε υφές και εξάγονται σε buffer εκτός οθόνης και οι αλγόριθμοι πολλαπλών περασμάτων χρησιμοποιούν μονάδες σκίασης pixel. Οι περιορισμοί GPGPU μπορεί να περιλαμβάνουν: ανεπαρκή χρήση δυνατοτήτων υλικού, περιορισμούς εύρους ζώνης μνήμης, έλλειψη λειτουργίας scatter (μόνο συλλογή), υποχρεωτική χρήση του API γραφικών.

Τα κύρια πλεονεκτήματα του CUDA έναντι των προηγούμενων μεθόδων GPGPU πηγάζουν από το γεγονός ότι η αρχιτεκτονική έχει σχεδιαστεί για να κάνει αποτελεσματική χρήση υπολογιστών μη γραφικών στη GPU και χρησιμοποιεί τη γλώσσα προγραμματισμού C χωρίς να απαιτείται η μεταφορά αλγορίθμων σε μια φιλική προς την ιδέα μορφή αγωγού γραφικών . Το CUDA προσφέρει μια νέα διαδρομή προς την υπολογιστική GPU που δεν χρησιμοποιεί API γραφικών, προσφέροντας τυχαία πρόσβαση στη μνήμη (διασκορπισμός ή συγκέντρωση). Αυτή η αρχιτεκτονική δεν έχει τα μειονεκτήματα της GPGPU και χρησιμοποιεί όλες τις μονάδες εκτέλεσης και επεκτείνει επίσης τις δυνατότητες λόγω μαθηματικών ακεραίων και λειτουργιών μετατόπισης bit.

Επιπλέον, το CUDA ανοίγει ορισμένες δυνατότητες υλικού που δεν είναι διαθέσιμες από τα API γραφικών, όπως η κοινόχρηστη μνήμη. Αυτή είναι μια μικρή μνήμη (16 kilobyte ανά πολυεπεξεργαστή) στην οποία έχουν πρόσβαση τα μπλοκ νημάτων. Σας επιτρέπει να αποθηκεύετε προσωρινά τα δεδομένα με τη μεγαλύτερη πρόσβαση και μπορεί να παρέχει μεγαλύτερες ταχύτητες από τη χρήση ανακτήσεων υφής για αυτήν την εργασία. Το οποίο, με τη σειρά του, μειώνει την ευαισθησία διεκπεραίωσης των παράλληλων αλγορίθμων σε πολλές εφαρμογές. Για παράδειγμα, είναι χρήσιμο για γραμμική άλγεβρα, γρήγορο μετασχηματισμό Fourier και φίλτρα επεξεργασίας εικόνας.

Η πρόσβαση στη μνήμη είναι επίσης πιο βολική στο CUDA. Ο κώδικας API γραφικών εξάγει δεδομένα ως 32 τιμές κινητής υποδιαστολής μονής ακρίβειας (τιμές RGBA ταυτόχρονα σε οκτώ στόχους απόδοσης) σε προκαθορισμένες περιοχές και το CUDA υποστηρίζει scatter εγγραφή - απεριόριστο αριθμό εγγραφών σε οποιαδήποτε διεύθυνση. Τέτοια πλεονεκτήματα καθιστούν δυνατή την εκτέλεση ορισμένων αλγορίθμων στη GPU που δεν μπορούν να εφαρμοστούν αποτελεσματικά χρησιμοποιώντας μεθόδους GPGPU που βασίζονται σε API γραφικών.

Επίσης, τα API γραφικών αποθηκεύουν απαραίτητα δεδομένα σε textures, κάτι που απαιτεί προκαταρκτική συσκευασία μεγάλων συστοιχιών σε textures, γεγονός που περιπλέκει τον αλγόριθμο και επιβάλλει τη χρήση ειδικών διευθύνσεων. Και το CUDA σάς επιτρέπει να διαβάζετε δεδομένα σε οποιαδήποτε διεύθυνση. Ένα άλλο πλεονέκτημα του CUDA είναι η βελτιστοποιημένη ανταλλαγή δεδομένων μεταξύ της CPU και της GPU. Και για προγραμματιστές που θέλουν πρόσβαση χαμηλού επιπέδου (για παράδειγμα, όταν γράφουν άλλη γλώσσα προγραμματισμού), το CUDA προσφέρει δυνατότητες προγραμματισμού γλώσσας assembly χαμηλού επιπέδου.

Ιστορία ανάπτυξης CUDA

Η ανάπτυξη του CUDA ανακοινώθηκε μαζί με το τσιπ G80 τον Νοέμβριο του 2006 και η κυκλοφορία μιας δημόσιας beta έκδοσης του CUDA SDK πραγματοποιήθηκε τον Φεβρουάριο του 2007. Η έκδοση 1.0 κυκλοφόρησε τον Ιούνιο του 2007 για να λανσάρει λύσεις Tesla που βασίζονται στο τσιπ G80 και προορίζονται για την αγορά υπολογιστών υψηλής απόδοσης. Στη συνέχεια, στα τέλη του έτους, κυκλοφόρησε η beta έκδοση του CUDA 1.1, η οποία, παρά τη μικρή αύξηση του αριθμού έκδοσης, εισήγαγε αρκετά νέα πράγματα.

Το νέο στο CUDA 1.1 είναι η συμπερίληψη της λειτουργικότητας CUDA στα κανονικά προγράμματα οδήγησης βίντεο της Nvidia. Αυτό σήμαινε ότι στις απαιτήσεις για οποιοδήποτε πρόγραμμα CUDA αρκούσε να υποδεικνύεται μια κάρτα βίντεο της σειράς Geforce 8 και άνω, καθώς και μια ελάχιστη έκδοση προγράμματος οδήγησης 169.xx. Αυτό είναι πολύ σημαντικό για τους προγραμματιστές, εάν πληρούνται αυτές οι προϋποθέσεις, τα προγράμματα CUDA θα λειτουργήσουν για οποιονδήποτε χρήστη. Προστέθηκαν επίσης η ασύγχρονη εκτέλεση μαζί με την αντιγραφή δεδομένων (μόνο για τσιπ G84, G86, G92 και υψηλότερα), η ασύγχρονη μεταφορά δεδομένων στη μνήμη βίντεο, οι λειτουργίες πρόσβασης στην ατομική μνήμη, η υποστήριξη για εκδόσεις 64 bit των Windows και η δυνατότητα λειτουργίας πολλαπλών chip CUDA σε λειτουργία SLI.

Προς το παρόν, η τρέχουσα έκδοση είναι για λύσεις που βασίζονται στο GT200 CUDA 2.0, που κυκλοφόρησε μαζί με τη σειρά Geforce GTX 200. Η έκδοση beta κυκλοφόρησε την άνοιξη του 2008. Προστέθηκε η δεύτερη έκδοση: υποστήριξη για υπολογισμούς διπλής ακρίβειας (υποστήριξη υλικού μόνο για το GT200), τέλος υποστηρίζει Windows Vista (εκδόσεις 32 και 64 bit) και Mac OS X, πρόσθετα εργαλεία εντοπισμού σφαλμάτων και προφίλ, υποστηρίζει υφές 3D, βελτιστοποιημένη μεταφορά δεδομένων .

Όσον αφορά τους υπολογισμούς διπλής ακρίβειας, η ταχύτητά τους στην τρέχουσα παραγωγή υλικού είναι αρκετές φορές χαμηλότερη από την απλή ακρίβεια. Οι λόγοι συζητούνται στα δικά μας. Η εφαρμογή αυτής της υποστήριξης στο GT200 είναι ότι τα μπλοκ FP32 δεν χρησιμοποιούνται για τη λήψη αποτελεσμάτων με τέσσερις φορές χαμηλότερη ταχύτητα για την υποστήριξη υπολογισμών FP64, η Nvidia αποφάσισε να δημιουργήσει αποκλειστικές υπολογιστικές μονάδες. Και στο GT200 υπάρχουν δέκα φορές λιγότερα από αυτά τα μπλοκ FP32 (ένα μπλοκ διπλής ακρίβειας για κάθε πολυεπεξεργαστή).

Στην πραγματικότητα, η απόδοση μπορεί να είναι ακόμη χαμηλότερη, καθώς η αρχιτεκτονική είναι βελτιστοποιημένη για ανάγνωση 32-bit από τη μνήμη και τους καταχωρητές, επιπλέον, δεν απαιτείται διπλή ακρίβεια σε εφαρμογές γραφικών και στο GT200 είναι πιο πιθανό να υπάρχει. Και οι σύγχρονοι τετραπύρηνες επεξεργαστές δείχνουν όχι πολύ λιγότερη πραγματική απόδοση. Ωστόσο, όντας ακόμη και 10 φορές πιο αργή από την απλή ακρίβεια, μια τέτοια υποστήριξη είναι χρήσιμη για σχέδια μικτής ακρίβειας. Μια κοινή τεχνική είναι να λαμβάνονται αρχικά κατά προσέγγιση αποτελέσματα με απλή ακρίβεια και στη συνέχεια να τα τελειοποιούνται με διπλή ακρίβεια. Τώρα αυτό μπορεί να γίνει απευθείας στην κάρτα βίντεο, χωρίς αποστολή ενδιάμεσων δεδομένων στην CPU.

Ένα άλλο χρήσιμο χαρακτηριστικό του CUDA 2.0 δεν έχει καμία σχέση με την GPU, παραδόξως. Απλώς, μπορείτε τώρα να μεταγλωττίσετε τον κώδικα CUDA σε πολύ αποδοτικό κώδικα SSE πολλαπλών νημάτων για γρήγορη εκτέλεση στην CPU. Δηλαδή, τώρα αυτή η δυνατότητα είναι κατάλληλη όχι μόνο για εντοπισμό σφαλμάτων, αλλά και για πραγματική χρήση σε συστήματα χωρίς κάρτα βίντεο Nvidia. Εξάλλου, η χρήση του CUDA σε κανονικό κώδικα παρεμποδίζεται από το γεγονός ότι οι κάρτες γραφικών Nvidia, αν και οι πιο δημοφιλείς μεταξύ αποκλειστικών λύσεων βίντεο, δεν είναι διαθέσιμες σε όλα τα συστήματα. Και πριν από την έκδοση 2.0, σε τέτοιες περιπτώσεις θα ήταν απαραίτητο να δημιουργηθούν δύο διαφορετικοί κωδικοί: για CUDA και χωριστά για την CPU. Και τώρα μπορείτε να εκτελέσετε οποιοδήποτε πρόγραμμα CUDA σε μια CPU με υψηλή απόδοση, αν και με χαμηλότερη ταχύτητα από ό,τι σε τσιπ βίντεο.

Λύσεις που υποστηρίζουν Nvidia CUDA

Όλες οι κάρτες γραφικών με δυνατότητα CUDA μπορούν να βοηθήσουν στην επιτάχυνση των πιο απαιτητικών εργασιών, από την επεξεργασία ήχου και βίντεο έως την ιατρική και την επιστημονική έρευνα. Ο μόνος πραγματικός περιορισμός είναι ότι πολλά προγράμματα CUDA απαιτούν τουλάχιστον 256 megabyte μνήμης βίντεο και αυτό είναι ένα από τα πιο σημαντικά τεχνικά χαρακτηριστικά για τις εφαρμογές CUDA.

Μπορείτε να βρείτε την πιο πρόσφατη λίστα προϊόντων που υποστηρίζουν CUDA στη διεύθυνση. Κατά τη στιγμή της γραφής, οι υπολογισμοί CUDA υποστήριζαν όλα τα προϊόντα των σειρών GeForce 200, GeForce 9 και GeForce 8, συμπεριλαμβανομένων των προϊόντων κινητής τηλεφωνίας που ξεκινούν με το GeForce 8400M, καθώς και τα chipset GeForce 8100, 8200 και 8300 Modern Quadro και 830 υποστήριξη CUDA όλων των Tesla: S1070, C1060, C870, D870 και S870.

Σημειώνουμε ιδιαίτερα ότι μαζί με τις νέες κάρτες γραφικών Geforce GTX 260 και 280, ανακοινώθηκαν και αντίστοιχες λύσεις για υπολογιστές υψηλής απόδοσης: Tesla C1060 και S1070 (φαίνεται στην παραπάνω φωτογραφία), οι οποίες θα είναι διαθέσιμες για αγορά αυτό το φθινόπωρο. Χρησιμοποιούν την ίδια GPU - GT200, στο C1060 υπάρχει ένα, στο S1070 υπάρχουν τέσσερις. Όμως, σε αντίθεση με τις λύσεις gaming, χρησιμοποιούν τέσσερα gigabyte μνήμης ανά τσιπ. Το μόνο μειονέκτημα είναι ότι η συχνότητα μνήμης και το εύρος ζώνης είναι χαμηλότερα από αυτά των καρτών gaming, παρέχοντας 102 GB/s ανά τσιπ.

Σύνθεση της Nvidia CUDA

Το CUDA περιλαμβάνει δύο API: υψηλού επιπέδου (CUDA Runtime API) και χαμηλού επιπέδου (CUDA Driver API), αν και είναι αδύνατο να χρησιμοποιήσετε και τα δύο ταυτόχρονα σε ένα πρόγραμμα, πρέπει να χρησιμοποιήσετε το ένα ή το άλλο. Το υψηλού επιπέδου λειτουργεί "πάνω" από το χαμηλότερου επιπέδου, όλες οι κλήσεις χρόνου εκτέλεσης μεταφράζονται σε απλές οδηγίες που επεξεργάζονται το χαμηλού επιπέδου Driver API. Αλλά ακόμη και ένα API "υψηλού επιπέδου" προϋποθέτει γνώση σχετικά με το σχεδιασμό και τη λειτουργία των τσιπ βίντεο της Nvidia, δεν υπάρχει πολύ υψηλό επίπεδο αφαίρεσης.

Υπάρχει ένα άλλο επίπεδο, ακόμη υψηλότερο - δύο βιβλιοθήκες:

CUBLASΈκδοση CUDA του BLAS (Βασικά Υποπρογράμματα Γραμμικής Άλγεβρας), σχεδιασμένη για υπολογισμό προβλημάτων γραμμικής άλγεβρας και χρήση άμεσης πρόσβασης σε πόρους GPU.

CUFFTΗ έκδοση CUDA της βιβλιοθήκης Fast Fourier Transform για τον υπολογισμό του γρήγορου μετασχηματισμού Fourier, που χρησιμοποιείται ευρέως στην επεξεργασία σήματος. Υποστηρίζονται οι ακόλουθοι τύποι μετασχηματισμού: σύμπλοκο-σύνθετο (C2C), πραγματικό-σύνθετο (R2C) και σύμπλοκο-πραγματικό (C2R).

Ας ρίξουμε μια πιο προσεκτική ματιά σε αυτές τις βιβλιοθήκες. Οι CUBLAS είναι τυπικοί αλγόριθμοι γραμμικής άλγεβρας μεταφρασμένοι στη γλώσσα CUDA. Η βιβλιοθήκη είναι πολύ εύκολη στη χρήση: πρέπει να δημιουργήσετε μια μήτρα και διανυσματικά αντικείμενα στη μνήμη της κάρτας βίντεο, να τα γεμίσετε με δεδομένα, να καλέσετε τις απαιτούμενες λειτουργίες CUBLAS και να φορτώσετε τα αποτελέσματα από τη μνήμη βίντεο πίσω στη μνήμη του συστήματος. Το CUBLAS περιέχει ειδικές λειτουργίες για τη δημιουργία και την καταστροφή αντικειμένων στη μνήμη GPU, καθώς και για την ανάγνωση και εγγραφή δεδομένων σε αυτήν τη μνήμη. Υποστηριζόμενες λειτουργίες BLAS: επίπεδα 1, 2 και 3 για πραγματικούς αριθμούς, επίπεδο CGEMM 1 για μιγαδικούς αριθμούς. Το επίπεδο 1 είναι πράξεις διανύσματος-διανύσματος, το επίπεδο 2 είναι πράξεις διανύσματος-μήτρας, το επίπεδο 3 είναι πράξεις μήτρας-μήτρας.

Η έκδοση CUFFT CUDA της συνάρτησης Fast Fourier Transform, ευρέως χρησιμοποιούμενη και πολύ σημαντική στην ανάλυση σήματος, το φιλτράρισμα κ.λπ. Το CUFFT παρέχει μια απλή διεπαφή για τον αποτελεσματικό υπολογισμό του FFT σε GPU της Nvidia χωρίς να χρειάζεται να αναπτύξετε τη δική σας GPU FFT. Η παραλλαγή CUDA του FFT υποστηρίζει μετασχηματισμούς 1D, 2D και 3D σύνθετων και πραγματικών δεδομένων, εκτέλεση δέσμης για πολλαπλούς μετασχηματισμούς 1D παράλληλα, τα μεγέθη των 2D και 3D μετασχηματισμών μπορεί να είναι εντός , για 1D το μέγεθος έως και 8 εκατομμυρίων στοιχείων είναι υποστηρίζεται.

Βασικά στοιχεία δημιουργίας προγραμμάτων στο CUDA

Για να κατανοήσετε περαιτέρω κείμενο, θα πρέπει να κατανοήσετε τα βασικά αρχιτεκτονικά χαρακτηριστικά των τσιπ βίντεο Nvidia. Η GPU αποτελείται από πολλά συμπλέγματα μονάδων υφής (Texture Processing Cluster). Κάθε σύμπλεγμα αποτελείται από ένα μεγάλο μπλοκ ανακτήσεων υφής και δύο ή τρεις πολυεπεξεργαστές ροής, καθένας από τους οποίους αποτελείται από οκτώ υπολογιστικές συσκευές και δύο υπερλειτουργικές μονάδες. Όλες οι εντολές εκτελούνται χρησιμοποιώντας την αρχή SIMD, όπου μια εντολή εφαρμόζεται σε όλα τα νήματα σε ένα στημόνι (όρος από την κλωστοϋφαντουργία, στο CUDA αυτό είναι μια ομάδα 32 νημάτων ο ελάχιστος όγκος δεδομένων που επεξεργάζονται οι πολυεπεξεργαστές). Αυτή η μέθοδος εκτέλεσης ονομαζόταν SIMT (single instruction multiple threads one instruction and many threads).

Κάθε ένας από τους πολυεπεξεργαστές έχει ορισμένους πόρους. Έτσι, υπάρχει μια ειδική κοινή μνήμη 16 kilobyte ανά πολυεπεξεργαστή. Αλλά αυτό δεν είναι προσωρινή μνήμη, αφού ο προγραμματιστής μπορεί να το χρησιμοποιήσει για οποιαδήποτε ανάγκη, όπως το Local Store στο SPU των επεξεργαστών Cell. Αυτή η κοινόχρηστη μνήμη επιτρέπει την ανταλλαγή πληροφοριών μεταξύ νημάτων του ίδιου μπλοκ. Είναι σημαντικό όλα τα νήματα ενός μπλοκ να εκτελούνται πάντα από τον ίδιο πολυεπεξεργαστή. Αλλά τα νήματα από διαφορετικά μπλοκ δεν μπορούν να ανταλλάξουν δεδομένα και πρέπει να θυμάστε αυτόν τον περιορισμό. Η κοινόχρηστη μνήμη είναι συχνά χρήσιμη εκτός από την περίπτωση που πολλά νήματα έχουν πρόσβαση στην ίδια τράπεζα μνήμης. Οι πολυεπεξεργαστές μπορούν επίσης να έχουν πρόσβαση στη μνήμη βίντεο, αλλά με υψηλότερες καθυστερήσεις και χειρότερη απόδοση. Για να επιταχύνουν την πρόσβαση και να μειώσουν τη συχνότητα πρόσβασης στη μνήμη βίντεο, οι πολυεπεξεργαστές διαθέτουν 8 kilobyte κρυφής μνήμης για δεδομένα σταθερών και υφής.

Ο πολυεπεξεργαστής χρησιμοποιεί καταχωρητές 8192-16384 (για G8x/G9x και GT2xx, αντίστοιχα), κοινά σε όλα τα νήματα όλων των μπλοκ που εκτελούνται σε αυτόν. Ο μέγιστος αριθμός μπλοκ ανά πολυεπεξεργαστή για το G8x/G9x είναι οκτώ και ο αριθμός των στημονιών είναι 24 (768 νήματα ανά πολυεπεξεργαστή). Συνολικά, οι κορυφαίες κάρτες γραφικών των σειρών Geforce 8 και 9 μπορούν να επεξεργαστούν έως και 12288 νήματα τη φορά. Η GeForce GTX 280 που βασίζεται στο GT200 προσφέρει έως και 1024 νήματα ανά πολυεπεξεργαστή, διαθέτει 10 συμπλέγματα τριών πολυεπεξεργαστών που επεξεργάζονται έως και 30720 νήματα. Η γνώση αυτών των περιορισμών σάς επιτρέπει να βελτιστοποιήσετε αλγόριθμους για τους διαθέσιμους πόρους.

Το πρώτο βήμα κατά τη μετεγκατάσταση μιας υπάρχουσας εφαρμογής στο CUDA είναι η δημιουργία προφίλ και ο εντοπισμός περιοχών κώδικα που αποτελούν σημεία συμφόρησης που επιβραδύνουν την εργασία. Εάν ανάμεσα σε τέτοια τμήματα υπάρχουν εκείνα που είναι κατάλληλα για γρήγορη παράλληλη εκτέλεση, αυτές οι συναρτήσεις μεταφέρονται σε επεκτάσεις CUDA C για εκτέλεση στη GPU. Το πρόγραμμα μεταγλωττίζεται χρησιμοποιώντας έναν μεταγλωττιστή που παρέχεται από την Nvidia που δημιουργεί κώδικα τόσο για την CPU όσο και για την GPU. Όταν εκτελείται ένα πρόγραμμα, ο κεντρικός επεξεργαστής εκτελεί τα τμήματα του κώδικα και η GPU εκτελεί τον κώδικα CUDA με τους περισσότερους παράλληλους υπολογισμούς. Αυτό το τμήμα αφιερωμένο στην GPU ονομάζεται πυρήνας. Ο πυρήνας ορίζει τις λειτουργίες που θα εκτελεστούν στα δεδομένα.

Το τσιπ βίντεο λαμβάνει τον πυρήνα και δημιουργεί αντίγραφα για κάθε στοιχείο δεδομένων. Αυτά τα αντίγραφα ονομάζονται νήματα. Μια ροή περιέχει έναν μετρητή, καταχωρητές και κατάσταση. Για μεγάλους όγκους δεδομένων, όπως η επεξεργασία εικόνας, εκκινούνται εκατομμύρια νήματα. Τα νήματα εκτελούνται σε ομάδες των 32, που ονομάζονται στρέβλωση, εκχωρούνται σε συγκεκριμένους πολυεπεξεργαστές νημάτων. Κάθε πολυεπεξεργαστής αποτελείται από οκτώ πυρήνες επεξεργαστές ροής που εκτελούν μία εντολή MAD ανά κύκλο ρολογιού. Για να εκτελέσετε ένα στρέβλωμα 32 νημάτων, απαιτούνται τέσσερις κύκλοι ρολογιού του πολυεπεξεργαστή (μιλάμε για τη συχνότητα του τομέα shader, η οποία είναι 1,5 GHz και άνω).

Ο πολυεπεξεργαστής δεν είναι ένας παραδοσιακός πολυπύρηνος επεξεργαστής, είναι εξαιρετικά ικανός για πολλαπλές νήματα, υποστηρίζοντας έως και 32 παραμορφώσεις κάθε φορά, το υλικό επιλέγει ποιο στρέβλωμα θα εκτελέσει και αλλάζει από το ένα στο άλλο. κύκλους ρολογιού. Εάν σχεδιάσουμε μια αναλογία με έναν κεντρικό επεξεργαστή, αυτό είναι παρόμοιο με την ταυτόχρονη εκτέλεση 32 προγραμμάτων και την εναλλαγή μεταξύ τους σε κάθε κύκλο ρολογιού χωρίς απώλεια της εναλλαγής περιβάλλοντος. Στην πραγματικότητα, οι πυρήνες της CPU υποστηρίζουν την ταυτόχρονη εκτέλεση ενός προγράμματος και μεταβαίνουν σε άλλα με καθυστέρηση εκατοντάδων κύκλων.

Μοντέλο προγραμματισμού CUDA

Ας επαναλάβουμε ότι το CUDA χρησιμοποιεί ένα παράλληλο υπολογιστικό μοντέλο, όταν κάθε ένας από τους επεξεργαστές SIMD εκτελεί την ίδια εντολή σε διαφορετικά στοιχεία δεδομένων παράλληλα. Η GPU είναι μια υπολογιστική συσκευή, ένας συνεπεξεργαστής (συσκευή) για τον κεντρικό επεξεργαστή (host), ο οποίος έχει τη δική του μνήμη και επεξεργάζεται μεγάλο αριθμό νημάτων παράλληλα. Ο πυρήνας είναι μια συνάρτηση για τη GPU, που εκτελείται από νήματα (μια αναλογία από τα τρισδιάστατα γραφικά - ένα shader).

Είπαμε παραπάνω ότι ένα τσιπ βίντεο διαφέρει από μια CPU στο ότι μπορεί να επεξεργαστεί δεκάδες χιλιάδες νήματα ταυτόχρονα, κάτι που είναι χαρακτηριστικό για γραφικά που είναι καλά παραλληλισμένα. Κάθε ροή είναι κλιμακωτή και δεν απαιτεί τη συσκευασία δεδομένων σε διανύσματα 4 συστατικών, κάτι που είναι πιο βολικό για τις περισσότερες εργασίες. Ο αριθμός των λογικών νημάτων και των μπλοκ νημάτων υπερβαίνει τον αριθμό των φυσικών συσκευών εκτέλεσης, γεγονός που παρέχει καλή επεκτασιμότητα για ολόκληρο το φάσμα των εταιρικών λύσεων.

Το μοντέλο προγραμματισμού CUDA περιλαμβάνει ομαδοποίηση νημάτων. Τα νήματα οργανώνονται σε μπλοκ νημάτων—μονοδιάστατα ή δισδιάστατα πλέγματα νημάτων που επικοινωνούν μεταξύ τους χρησιμοποιώντας κοινόχρηστη μνήμη και σημεία συγχρονισμού. Το πρόγραμμα (πυρήνας) τρέχει πάνω από ένα πλέγμα μπλοκ νημάτων, δείτε το παρακάτω σχήμα. Ένα πλέγμα εκτελείται κάθε φορά. Κάθε μπλοκ μπορεί να είναι μονοδιάστατο, δύο ή τρισδιάστατο σε σχήμα και μπορεί να αποτελείται από 512 νήματα στο τρέχον υλικό.

Τα μπλοκ νημάτων εκτελούνται σε μικρές ομάδες που ονομάζονται στημόνι, το μέγεθος των οποίων είναι 32 νήματα. Αυτή είναι η ελάχιστη ποσότητα δεδομένων που μπορεί να υποβληθεί σε επεξεργασία σε πολυεπεξεργαστές. Και επειδή αυτό δεν είναι πάντα βολικό, το CUDA σάς επιτρέπει να εργάζεστε με μπλοκ που περιέχουν από 64 έως 512 νήματα.

Η ομαδοποίηση μπλοκ σε πλέγματα σάς επιτρέπει να ξεφύγετε από τους περιορισμούς και να εφαρμόσετε τον πυρήνα σε περισσότερα νήματα σε μία μόνο κλήση. Αυτό βοηθά επίσης στην απολέπιση. Εάν η GPU δεν έχει αρκετούς πόρους, θα εκτελεί μπλοκ διαδοχικά. Διαφορετικά, τα μπλοκ μπορούν να εκτελεστούν παράλληλα, κάτι που είναι σημαντικό για τη βέλτιστη κατανομή της εργασίας σε τσιπ βίντεο διαφορετικών επιπέδων, που κυμαίνονται από κινητά και ενσωματωμένα.

Μοντέλο μνήμης CUDA

Το μοντέλο μνήμης στο CUDA διακρίνεται από τη δυνατότητα διευθυνσιοδότησης byte, υποστήριξη τόσο για συλλογή όσο και για διασπορά. Ένας αρκετά μεγάλος αριθμός καταχωρητών είναι διαθέσιμος για κάθε επεξεργαστή ροής, έως και 1024 τεμάχια. Η πρόσβαση σε αυτά είναι πολύ γρήγορη και μπορείτε να αποθηκεύσετε ακέραιους αριθμούς 32 bit ή αριθμούς κινητής υποδιαστολής σε αυτούς.

Κάθε νήμα έχει πρόσβαση στους ακόλουθους τύπους μνήμης:

Καθολική μνήμηη μεγαλύτερη διαθέσιμη μνήμη για όλους τους πολυεπεξεργαστές σε ένα τσιπ βίντεο, το μέγεθος κυμαίνεται από 256 megabyte έως 1,5 gigabyte στις τρέχουσες λύσεις (και έως 4 GB στο Tesla). Έχει υψηλή απόδοση, περισσότερα από 100 gigabyte/s για κορυφαίες λύσεις Nvidia, αλλά πολύ υψηλές καθυστερήσεις αρκετών εκατοντάδων κύκλων. Δεν υπάρχει προσωρινή αποθήκευση, υποστηρίζει γενικές οδηγίες φόρτωσης και αποθήκευσης και δείκτες κανονικής μνήμης.

Τοπική μνήμηΑυτή είναι μια μικρή ποσότητα μνήμης στην οποία έχει πρόσβαση μόνο ένας επεξεργαστής ροής. Είναι σχετικά αργό - το ίδιο με το παγκόσμιο.

Κοινή μνήμηΑυτό είναι ένα μπλοκ μνήμης 16 kilobyte (σε τσιπ βίντεο της τρέχουσας αρχιτεκτονικής) με κοινόχρηστη πρόσβαση για όλους τους επεξεργαστές ροής στον πολυεπεξεργαστή. Αυτή η μνήμη είναι πολύ γρήγορη, όπως και οι καταχωρητές. Επιτρέπει στα νήματα να επικοινωνούν, ελέγχεται άμεσα από τον προγραμματιστή και έχει χαμηλή καθυστέρηση. Πλεονεκτήματα της κοινόχρηστης μνήμης: χρήση ως μνήμη cache πρώτου επιπέδου ελεγχόμενη από προγραμματιστή, μειωμένη καθυστέρηση κατά την πρόσβαση σε μονάδες εκτέλεσης (ALU) σε δεδομένα, μειωμένος αριθμός προσβάσεων στην καθολική μνήμη.

Σταθερή μνήμη- περιοχή μνήμης 64 kilobyte (το ίδιο για τις τρέχουσες GPU), μόνο για ανάγνωση από όλους τους πολυεπεξεργαστές. Αποθηκεύεται προσωρινά στα 8 kilobyte ανά πολυεπεξεργαστή. Αρκετά αργή - καθυστέρηση αρκετών εκατοντάδων κύκλων ελλείψει των απαραίτητων δεδομένων στη μνήμη cache.

Μνήμη υφήςένα μπλοκ μνήμης που μπορεί να διαβαστεί από όλους τους πολυεπεξεργαστές. Η δειγματοληψία δεδομένων πραγματοποιείται χρησιμοποιώντας τα μπλοκ υφής του τσιπ βίντεο, επομένως παρέχεται γραμμική παρεμβολή δεδομένων χωρίς πρόσθετο κόστος. 8 kilobyte αποθηκεύονται στην κρυφή μνήμη ανά πολυεπεξεργαστή. Αργή ως παγκόσμιοι εκατοντάδες κύκλοι λανθάνοντος χρόνου όταν δεν υπάρχουν δεδομένα στην κρυφή μνήμη.

Φυσικά, η συνολική, η τοπική, η υφή και η σταθερή μνήμη είναι φυσικά η ίδια μνήμη, γνωστή ως τοπική μνήμη βίντεο της κάρτας βίντεο. Οι διαφορές τους έγκεινται σε διαφορετικούς αλγόριθμους προσωρινής αποθήκευσης και μοντέλα πρόσβασης. Η CPU μπορεί να ενημερώσει και να ζητήσει μόνο εξωτερική μνήμη: καθολική, σταθερή και υφή.

Από όσα γράφτηκαν παραπάνω, είναι σαφές ότι το CUDA υιοθετεί μια ειδική προσέγγιση στην ανάπτυξη, όχι ακριβώς ίδια με αυτή που υιοθετείται σε προγράμματα για τη CPU. Πρέπει να θυμάστε για διαφορετικούς τύπους μνήμης, ότι η τοπική και η καθολική μνήμη δεν αποθηκεύονται στην κρυφή μνήμη και ότι ο λανθάνοντας χρόνος πρόσβασης σε αυτήν είναι πολύ υψηλότερος από αυτόν της μνήμης εγγραφής, καθώς βρίσκεται φυσικά σε ξεχωριστά τσιπ.

Ένα τυπικό, αλλά όχι απαραίτητο, μοτίβο επίλυσης προβλημάτων:

• η εργασία χωρίζεται σε δευτερεύουσες εργασίες.
• Τα δεδομένα εισόδου χωρίζονται σε μπλοκ που χωρούν στην κοινόχρηστη μνήμη.
• κάθε μπλοκ επεξεργάζεται από ένα μπλοκ νημάτων.
• το υπομπλοκ φορτώνεται στην κοινόχρηστη μνήμη από την καθολική μνήμη.
• διενεργούνται κατάλληλοι υπολογισμοί στα δεδομένα στην κοινόχρηστη μνήμη.
• τα αποτελέσματα αντιγράφονται από την κοινόχρηστη μνήμη πίσω στην καθολική μνήμη.

Προγραμματιστικό περιβάλλον

Το CUDA περιλαμβάνει βιβλιοθήκες χρόνου εκτέλεσης:

• ένα κοινό τμήμα που παρέχει ενσωματωμένους τύπους διανυσμάτων και υποσύνολα κλήσεων RTL που υποστηρίζονται σε CPU και GPU.
• Στοιχείο CPU για τον έλεγχο μιας ή περισσότερων GPU.
• Ένα στοιχείο GPU που παρέχει λειτουργικότητα συγκεκριμένης GPU.

Η κύρια διαδικασία εφαρμογής CUDA εκτελείται σε έναν γενικό επεξεργαστή (host), εκτελεί πολλαπλά αντίγραφα διεργασιών πυρήνα στην κάρτα βίντεο. Ο κώδικας για την CPU κάνει τα εξής: αρχικοποιεί τη GPU, εκχωρεί μνήμη στην κάρτα βίντεο και το σύστημα, αντιγράφει σταθερές στη μνήμη της κάρτας βίντεο, εκκινεί πολλά αντίγραφα διεργασιών πυρήνα στην κάρτα βίντεο, αντιγράφει το αποτέλεσμα από τη μνήμη βίντεο, ελευθερώνει το μνήμη και εξόδους.

Ως παράδειγμα κατανόησης, εδώ είναι ο κώδικας CPU για την προσθήκη διανύσματος που παρουσιάζεται στο CUDA:

Οι συναρτήσεις που εκτελούνται από το τσιπ βίντεο έχουν τους ακόλουθους περιορισμούς: δεν υπάρχει αναδρομή, δεν υπάρχουν στατικές μεταβλητές μέσα στις συναρτήσεις και δεν υπάρχει μεταβλητός αριθμός ορισμάτων. Υποστηρίζονται δύο τύποι διαχείρισης μνήμης: γραμμική μνήμη με πρόσβαση μέσω δεικτών 32-bit και πίνακες CUDA με πρόσβαση μόνο μέσω λειτουργιών ανάκτησης υφής.

Τα προγράμματα CUDA μπορούν να αλληλεπιδράσουν με τα API γραφικών: για απόδοση δεδομένων που δημιουργούνται στο πρόγραμμα, για ανάγνωση αποτελεσμάτων απόδοσης και επεξεργασία τους χρησιμοποιώντας εργαλεία CUDA (για παράδειγμα, κατά την εφαρμογή φίλτρων μετα-επεξεργασίας). Για να επιτευχθεί αυτό, οι πόροι API γραφικών μπορούν να αντιστοιχιστούν (λαμβάνοντας τη διεύθυνση του πόρου) στον παγκόσμιο χώρο μνήμης CUDA. Υποστηρίζονται οι ακόλουθοι τύποι πόρων API γραφικών: Αντικείμενα buffer (PBO/VBO) σε OpenGL, buffer και υφές κορυφής (2D, 3D και χάρτες κυβισμού) Direct3D9.

Στάδια σύνταξης μιας εφαρμογής CUDA:

Τα αρχεία πηγαίου κώδικα CUDA C μεταγλωττίζονται χρησιμοποιώντας το πρόγραμμα NVCC, το οποίο είναι ένα περιτύλιγμα πάνω από άλλα εργαλεία και τα ονομάζει: cudacc, g++, cl, κ.λπ. C και κωδικός αντικειμένου PTX για το τσιπ βίντεο. Τα εκτελέσιμα αρχεία με κωδικό CUDA απαιτούν απαραίτητα τη βιβλιοθήκη χρόνου εκτέλεσης CUDA (cudart) και τη βιβλιοθήκη πυρήνα CUDA (cuda).

Βελτιστοποίηση προγραμμάτων στο CUDA

Φυσικά, είναι αδύνατο να εξεταστούν σοβαρά ζητήματα βελτιστοποίησης στον προγραμματισμό CUDA στο πλαίσιο ενός άρθρου ανασκόπησης. Επομένως, θα μιλήσουμε εν συντομία για τα βασικά πράγματα. Για να χρησιμοποιήσετε αποτελεσματικά τις δυνατότητες του CUDA, πρέπει να ξεχάσετε τις συνήθεις μεθόδους σύνταξης προγραμμάτων για την CPU και να χρησιμοποιήσετε αυτούς τους αλγόριθμους που είναι καλά παραλληλισμένοι σε χιλιάδες νήματα. Είναι επίσης σημαντικό να βρείτε το βέλτιστο μέρος για την αποθήκευση δεδομένων (μητρώα, κοινόχρηστη μνήμη κ.λπ.), να ελαχιστοποιήσετε τη μεταφορά δεδομένων μεταξύ της CPU και της GPU και να χρησιμοποιήσετε buffering.

Γενικά, κατά τη βελτιστοποίηση ενός προγράμματος CUDA, θα πρέπει να προσπαθήσετε να επιτύχετε τη βέλτιστη ισορροπία μεταξύ του μεγέθους και του αριθμού των μπλοκ. Περισσότερα νήματα ανά μπλοκ θα μειώσουν τον αντίκτυπο της καθυστέρησης της μνήμης, αλλά θα μειώσουν επίσης τον αριθμό των διαθέσιμων καταχωρητών. Επιπλέον, ένα μπλοκ 512 νημάτων είναι αναποτελεσματικό.

Μεταξύ των κύριων σημείων βελτιστοποίησης των προγραμμάτων CUDA: η πιο ενεργή χρήση της κοινόχρηστης μνήμης, καθώς είναι πολύ ταχύτερη από την παγκόσμια μνήμη βίντεο της κάρτας βίντεο. Η ανάγνωση και η εγγραφή από την παγκόσμια μνήμη θα πρέπει να συνενώνονται όποτε είναι δυνατόν. Για να γίνει αυτό, πρέπει να χρησιμοποιήσετε ειδικούς τύπους δεδομένων για την ανάγνωση και εγγραφή 32/64/128 bit δεδομένων ταυτόχρονα σε μία λειτουργία. Εάν οι αναγνώσεις είναι δύσκολο να συνδυαστούν, μπορείτε να δοκιμάσετε να χρησιμοποιήσετε ανακτήσεις υφής.

συμπεράσματα

Η αρχιτεκτονική λογισμικού και υλικού που παρουσιάζει η Nvidia για υπολογισμούς σε τσιπ βίντεο CUDA είναι κατάλληλη για την επίλυση ενός ευρέος φάσματος εξαιρετικά παράλληλων εργασιών. Το CUDA τρέχει σε ένα ευρύ φάσμα GPU της Nvidia και βελτιώνει το μοντέλο προγραμματισμού της GPU απλοποιώντας το και προσθέτοντας μεγάλο αριθμό λειτουργιών όπως κοινή μνήμη, συγχρονισμό νημάτων, υπολογισμούς διπλής ακρίβειας και λειτουργίες ακεραίων.

Το CUDA είναι μια τεχνολογία διαθέσιμη σε κάθε προγραμματιστή λογισμικού και μπορεί να χρησιμοποιηθεί από οποιονδήποτε προγραμματιστή που γνωρίζει τη γλώσσα C. Απλώς πρέπει να συνηθίσετε στο διαφορετικό πρότυπο προγραμματισμού που είναι εγγενές στον παράλληλο υπολογισμό. Αλλά εάν ο αλγόριθμος είναι, κατ' αρχήν, καλά παραλληλισμένος, τότε η μελέτη και ο χρόνος που αφιερώθηκε στον προγραμματισμό στο CUDA θα επιστρέψουν πολλές φορές.

Είναι πιθανό ότι λόγω της ευρείας χρήσης καρτών βίντεο στον κόσμο, η ανάπτυξη παράλληλων υπολογιστών σε GPU θα επηρεάσει σε μεγάλο βαθμό τη βιομηχανία υπολογιστών υψηλής απόδοσης. Οι δυνατότητες αυτές έχουν ήδη προκαλέσει μεγάλο ενδιαφέρον στους επιστημονικούς κύκλους και όχι μόνο σε αυτούς. Εξάλλου, οι πιθανές ευκαιρίες για επιτάχυνση αλγορίθμων που μπορούν εύκολα να παραλληλιστούν (σε διαθέσιμο υλικό, κάτι που δεν είναι λιγότερο σημαντικό) κατά δεκάδες φορές δεν είναι τόσο συνηθισμένες.

Οι επεξεργαστές γενικής χρήσης αναπτύσσονται μάλλον αργά και δεν έχουν τέτοια άλματα απόδοσης. Ουσιαστικά, ενώ ακούγεται σαν πολλά χρήματα, όποιος χρειάζεται γρήγορους υπολογιστές μπορεί τώρα να έχει έναν φθηνό προσωπικό υπερυπολογιστή στο γραφείο του, μερικές φορές χωρίς καν να επενδύσει επιπλέον, καθώς οι κάρτες γραφικών Nvidia είναι τόσο ευρέως διαθέσιμες. Για να μην αναφέρουμε την αυξημένη απόδοση όσον αφορά τα GFLOPS/$ και GFLOPS/W που αγαπούν τόσο πολύ οι κατασκευαστές GPU.

Το μέλλον πολλών υπολογιστών βρίσκεται ξεκάθαρα σε παράλληλους αλγόριθμους και σχεδόν όλες οι νέες λύσεις και πρωτοβουλίες κατευθύνονται προς αυτή την κατεύθυνση. Μέχρι στιγμής, ωστόσο, η ανάπτυξη νέων παραδειγμάτων βρίσκεται στο αρχικό στάδιο, πρέπει να δημιουργήσετε χειροκίνητα νήματα και να προγραμματίσετε την πρόσβαση στη μνήμη, γεγονός που περιπλέκει τις εργασίες σε σύγκριση με τον συμβατικό προγραμματισμό. Όμως η τεχνολογία CUDA έχει κάνει ένα βήμα προς τη σωστή κατεύθυνση και φαίνεται ξεκάθαρα μια επιτυχημένη λύση, ειδικά αν η Nvidia καταφέρει να πείσει όσο το δυνατόν περισσότερους προγραμματιστές για τα οφέλη και τις προοπτικές της.

Αλλά, φυσικά, οι GPU δεν θα αντικαταστήσουν τις CPU. Στη σημερινή τους μορφή δεν προορίζονται για αυτό. Τώρα που τα τσιπ βίντεο κινούνται σταδιακά προς την CPU, γίνονται όλο και πιο καθολικά (υπολογισμοί κινητής υποδιαστολής απλής και διπλής ακρίβειας, υπολογισμοί ακεραίων), οι CPU γίνονται όλο και πιο «παράλληλες», αποκτώντας μεγάλο αριθμό πυρήνων, τεχνολογίες πολλαπλών νημάτων , για να μην αναφέρουμε την εμφάνιση μπλοκ SIMD και ετερογενών έργων επεξεργαστή. Πιθανότατα, η GPU και η CPU απλώς θα συγχωνευθούν στο μέλλον. Είναι γνωστό ότι πολλές εταιρείες, μεταξύ των οποίων η Intel και η AMD, εργάζονται σε παρόμοια έργα. Και δεν έχει σημασία αν οι GPU απορροφώνται από την CPU ή το αντίστροφο.

Στο άρθρο, μιλήσαμε κυρίως για τα οφέλη του CUDA. Υπάρχει όμως και μια μύγα στην αλοιφή. Ένα από τα λίγα μειονεκτήματα του CUDA είναι η κακή φορητότητά του. Αυτή η αρχιτεκτονική λειτουργεί μόνο σε τσιπ βίντεο αυτής της εταιρείας, και όχι σε όλα, αλλά ξεκινώντας από τις σειρές Geforce 8 και 9 και τα αντίστοιχα Quadro και Tesla. Ναι, υπάρχουν πολλές τέτοιες λύσεις στον κόσμο. Αυτό είναι απλά υπέροχο, αλλά οι ανταγωνιστές προσφέρουν τις δικές τους λύσεις που διαφέρουν από το CUDA. Έτσι, η AMD έχει Stream Computing, η Intel θα έχει Ct στο μέλλον.

Ποια τεχνολογία θα κερδίσει, θα διαδοθεί και θα ζήσει περισσότερο από άλλες - μόνο ο χρόνος θα δείξει. Αλλά το CUDA έχει μια καλή ευκαιρία, καθώς σε σύγκριση με το Stream Computing, για παράδειγμα, αντιπροσωπεύει ένα πιο ανεπτυγμένο και εύχρηστο περιβάλλον προγραμματισμού στην κανονική γλώσσα C. Ίσως κάποιος τρίτος μπορεί να βοηθήσει στον προσδιορισμό εκδίδοντας κάποια γενική λύση. Για παράδειγμα, στην επόμενη ενημέρωση του DirectX στην έκδοση 11, η Microsoft υποσχέθηκε υπολογιστικά shaders, τα οποία θα μπορούσαν να γίνουν κάποιο είδος μέσης λύσης που ταιριάζει σε όλους ή σχεδόν σε όλους.

Κρίνοντας από τα προκαταρκτικά δεδομένα, αυτός ο νέος τύπος shader δανείζεται πολλά από το μοντέλο CUDA. Και προγραμματίζοντας σε αυτό το περιβάλλον τώρα, μπορείτε να αποκτήσετε άμεσα οφέλη και τις απαραίτητες δεξιότητες για το μέλλον. Από υπολογιστική σκοπιά υψηλής απόδοσης, το DirectX έχει επίσης το ευδιάκριτο μειονέκτημα της κακής φορητότητας, καθώς το API περιορίζεται στην πλατφόρμα των Windows. Ωστόσο, αναπτύσσεται ένα άλλο πρότυπο - η ανοιχτή πρωτοβουλία πολλαπλών πλατφορμών OpenCL, η οποία υποστηρίζεται από τις περισσότερες εταιρείες, συμπεριλαμβανομένων των Nvidia, AMD, Intel, IBM και πολλών άλλων.

Μην ξεχνάτε ότι το επόμενο άρθρο για το CUDA θα διερευνήσει συγκεκριμένες πρακτικές εφαρμογές επιστημονικών και άλλων μη γραφικών υπολογιστών που εκτελούνται από προγραμματιστές από διαφορετικά μέρη του πλανήτη μας χρησιμοποιώντας το Nvidia CUDA.

Στην ανάπτυξη των σύγχρονων επεξεργαστών, υπάρχει μια τάση προς μια σταδιακή αύξηση του αριθμού των πυρήνων, γεγονός που αυξάνει τις δυνατότητές τους στον παράλληλο υπολογισμό. Ωστόσο, εδώ και πολύ καιρό υπάρχουν GPU που είναι σημαντικά ανώτερες από τις CPU από αυτή την άποψη. Και αυτές οι δυνατότητες των GPU έχουν ήδη ληφθεί υπόψη από ορισμένες εταιρείες. Οι πρώτες απόπειρες χρήσης επιταχυντών γραφικών για υπολογισμούς χωρίς στόχο έχουν γίνει από τα τέλη της δεκαετίας του '90. Αλλά μόνο η εμφάνιση των shaders έγινε η ώθηση για την ανάπτυξη μιας εντελώς νέας τεχνολογίας και το 2003 εμφανίστηκε η έννοια του GPGPU (General-purpose graphics processing units). Σημαντικό ρόλο στην ανάπτυξη αυτής της πρωτοβουλίας έπαιξε η BrookGPU, η οποία είναι μια ειδική επέκταση για τη γλώσσα C Πριν από την εμφάνιση της BrookGPU, οι προγραμματιστές μπορούσαν να εργαστούν με GPU μόνο μέσω του Direct3D ή του OpenGL API. Ο Brook επέτρεψε στους προγραμματιστές να εργαστούν με ένα οικείο περιβάλλον και ο ίδιος ο μεταγλωττιστής, χρησιμοποιώντας ειδικές βιβλιοθήκες, υλοποίησε την αλληλεπίδραση με τη GPU σε χαμηλό επίπεδο.

Μια τέτοια πρόοδος δεν μπορούσε παρά να προσελκύσει την προσοχή των ηγετών αυτού του κλάδου - της AMD και της NVIDIA, οι οποίοι άρχισαν να αναπτύσσουν τις δικές τους πλατφόρμες λογισμικού για μη γραφικούς υπολογιστές στις κάρτες γραφικών τους. Κανείς δεν γνωρίζει καλύτερα από τους προγραμματιστές GPU όλες τις αποχρώσεις και τα χαρακτηριστικά των προϊόντων τους, γεγονός που επιτρέπει σε αυτές τις ίδιες εταιρείες να βελτιστοποιούν το πακέτο λογισμικού για συγκεκριμένες λύσεις υλικού όσο το δυνατόν πιο αποτελεσματικά. Επί του παρόντος, η NVIDIA αναπτύσσει την πλατφόρμα CUDA (Compute Unified Device Architecture) που καλεί μια παρόμοια τεχνολογία CTM (Close To Metal) ή AMD Stream Computing. Θα δούμε μερικές από τις δυνατότητες του CUDA και θα αξιολογήσουμε στην πράξη τις υπολογιστικές δυνατότητες του τσιπ γραφικών G92 της κάρτας γραφικών GeForce 8800 GT.

Αλλά πρώτα, ας δούμε μερικές από τις αποχρώσεις της εκτέλεσης υπολογισμών με χρήση GPU. Το κύριο πλεονέκτημά τους είναι ότι το τσιπ γραφικών σχεδιάστηκε αρχικά για να εκτελεί πολλαπλά νήματα, ενώ κάθε πυρήνας μιας συμβατικής CPU εκτελεί μια ροή διαδοχικών εντολών. Οποιαδήποτε σύγχρονη GPU είναι ένας πολυεπεξεργαστής που αποτελείται από πολλά συμπλέγματα υπολογιστών, με πολλές ALU σε καθεμία. Το πιο ισχυρό σύγχρονο τσιπ GT200 αποτελείται από 10 τέτοια συμπλέγματα, καθένα από τα οποία έχει 24 επεξεργαστές ροής. Η δοκιμασμένη κάρτα γραφικών GeForce 8800 GT που βασίζεται στο τσιπ G92 διαθέτει επτά μεγάλες υπολογιστικές μονάδες με 16 επεξεργαστές ροής η καθεμία. Οι CPU χρησιμοποιούν μπλοκ SIMD SSE για διανυσματικούς υπολογισμούς (πολλαπλά δεδομένα μιας εντολής - μία εντολή εκτελείται σε πολλαπλά δεδομένα), κάτι που απαιτεί μετατροπή των δεδομένων σε 4 διανύσματα. Η GPU επεξεργάζεται τα νήματα κλιμακωτά, π.χ. μία εντολή εφαρμόζεται σε πολλά νήματα (SIMT - πολλαπλά νήματα απλής εντολής). Αυτό εξοικονομεί τους προγραμματιστές από τη μετατροπή δεδομένων σε διανύσματα και επιτρέπει την αυθαίρετη διακλάδωση σε ροές. Κάθε υπολογιστική μονάδα GPU έχει άμεση πρόσβαση στη μνήμη. Και το εύρος ζώνης της μνήμης βίντεο είναι υψηλότερο, χάρη στη χρήση πολλών ξεχωριστών ελεγκτών μνήμης (στο κορυφαίο G200 υπάρχουν 8 κανάλια 64 bit) και υψηλές συχνότητες λειτουργίας.

Γενικά, σε ορισμένες εργασίες όταν εργάζεστε με μεγάλες ποσότητες δεδομένων, οι GPU είναι πολύ πιο γρήγορες από τις CPU. Παρακάτω βλέπετε μια απεικόνιση αυτής της δήλωσης:

Το γράφημα δείχνει τη δυναμική της αύξησης της απόδοσης της CPU και της GPU από το 2003. Η NVIDIA θέλει να αναφέρει αυτά τα δεδομένα ως διαφήμιση στα έγγραφά της, αλλά είναι μόνο θεωρητικοί υπολογισμοί και στην πραγματικότητα το χάσμα, φυσικά, μπορεί να αποδειχθεί πολύ μικρότερο.

Ωστόσο, όπως και να έχει, υπάρχει ένα τεράστιο δυναμικό GPU που μπορούν να χρησιμοποιηθούν και που απαιτεί μια συγκεκριμένη προσέγγιση στην ανάπτυξη λογισμικού. Όλα αυτά υλοποιούνται στο περιβάλλον υλικού και λογισμικού CUDA, το οποίο αποτελείται από πολλά επίπεδα λογισμικού - το υψηλού επιπέδου CUDA Runtime API και το χαμηλού επιπέδου CUDA Driver API.

Το CUDA χρησιμοποιεί την τυπική γλώσσα C για προγραμματισμό, η οποία είναι ένα από τα κύρια πλεονεκτήματά του για τους προγραμματιστές. Αρχικά, το CUDA περιλαμβάνει τις βιβλιοθήκες BLAS (βασικό πακέτο γραμμικής άλγεβρας) και FFT (μετασχηματισμός Fourier). Το CUDA έχει επίσης τη δυνατότητα αλληλεπίδρασης με τα API γραφικών OpenGL ή DirectX, τη δυνατότητα ανάπτυξης σε χαμηλό επίπεδο και χαρακτηρίζεται από μια βελτιστοποιημένη κατανομή των ροών δεδομένων μεταξύ της CPU και της GPU. Οι υπολογισμοί CUDA εκτελούνται ταυτόχρονα με τους γραφικούς, σε αντίθεση με την παρόμοια πλατφόρμα AMD, όπου εκκινείται μια ειδική εικονική μηχανή για υπολογισμούς στη GPU. Αλλά μια τέτοια «συγκατοίκηση» είναι επίσης γεμάτη σφάλματα εάν δημιουργείται μεγάλο φορτίο από το API γραφικών ενώ το CUDA εκτελείται ταυτόχρονα - εξάλλου, οι λειτουργίες γραφικών εξακολουθούν να έχουν μεγαλύτερη προτεραιότητα. Η πλατφόρμα είναι συμβατή με λειτουργικά συστήματα 32 και 64 bit Windows XP, Windows Vista, MacOS X και διάφορες εκδόσεις Linux. Η πλατφόρμα είναι ανοιχτή και στον ιστότοπο, εκτός από τα ειδικά προγράμματα οδήγησης για την κάρτα βίντεο, μπορείτε να κατεβάσετε πακέτα λογισμικού CUDA Toolkit, CUDA Developer SDK, συμπεριλαμβανομένου ενός μεταγλωττιστή, προγράμματος εντοπισμού σφαλμάτων, τυπικών βιβλιοθηκών και τεκμηρίωσης.

Όσον αφορά την πρακτική εφαρμογή του CUDA, για μεγάλο χρονικό διάστημα αυτή η τεχνολογία χρησιμοποιήθηκε μόνο για εξαιρετικά εξειδικευμένους μαθηματικούς υπολογισμούς στον τομέα της σωματιδιακής φυσικής, της αστροφυσικής, της ιατρικής ή της πρόβλεψης αλλαγών στη χρηματοπιστωτική αγορά κ.λπ. Αλλά αυτή η τεχνολογία γίνεται σταδιακά πιο κοντά στους απλούς χρήστες, ειδικότερα, εμφανίζονται ειδικά πρόσθετα για το Photoshop που μπορούν να χρησιμοποιήσουν την υπολογιστική ισχύ της GPU. Σε μια ειδική σελίδα μπορείτε να μελετήσετε ολόκληρη τη λίστα των προγραμμάτων που χρησιμοποιούν τις δυνατότητες του NVIDIA CUDA.

Ως πρακτική δοκιμή της νέας τεχνολογίας στην κάρτα γραφικών MSI NX8800GT-T2D256E-OC, θα χρησιμοποιήσουμε το πρόγραμμα TMPGEnc. Αυτό το προϊόν είναι εμπορικό (η πλήρης έκδοση κοστίζει 100 $), αλλά για κάρτες γραφικών MSI έρχεται ως μπόνους σε δοκιμαστική έκδοση για περίοδο 30 ημερών. Μπορείτε να κάνετε λήψη αυτής της έκδοσης από τον ιστότοπο του προγραμματιστή, αλλά για να εγκαταστήσετε το TMPGEnc 4.0 XPress MSI Special Edition χρειάζεστε τον αρχικό δίσκο με προγράμματα οδήγησης από την κάρτα MSI - χωρίς αυτόν το πρόγραμμα δεν θα εγκατασταθεί.

Για να εμφανίσετε τις πληρέστερες πληροφορίες σχετικά με τις υπολογιστικές δυνατότητες στο CUDA και να τις συγκρίνετε με άλλους προσαρμογείς βίντεο, μπορείτε να χρησιμοποιήσετε το ειδικό βοηθητικό πρόγραμμα CUDA-Z. Αυτές είναι οι πληροφορίες που δίνει για την κάρτα γραφικών GeForce 8800GT:

Σε σύγκριση με τα μοντέλα αναφοράς, το αντίγραφό μας λειτουργεί σε υψηλότερες συχνότητες: ο τομέας ράστερ είναι 63 MHz υψηλότερος από τον ονομαστικό και οι μονάδες σκίασης είναι ταχύτερες κατά 174 MHz και η μνήμη είναι 100 MHz ταχύτερη.

Θα συγκρίνουμε την ταχύτητα μετατροπής του ίδιου βίντεο HD κατά τον υπολογισμό μόνο με χρήση της CPU και με την πρόσθετη ενεργοποίηση του CUDA στο πρόγραμμα TMPGEnc στην ακόλουθη διαμόρφωση:

• Επεξεργαστής: Pentium Dual-Core E5200 2,5 GHz;
• Μητρική πλακέτα: Gigabyte P35-S3;
• Μνήμη: 2x1GB GoodRam PC6400 (5-5-5-18-2T)
• Κάρτα βίντεο: MSI NX8800GT-T2D256E-OC;
• Σκληρός δίσκος: 320 GB WD3200AAKS;
• Τροφοδοτικό: CoolerMaster eXtreme Power 500-PCAP;
• Λειτουργικό σύστημα: Windows XP SP2;
• TMPGEnc 4.0 XPress 4.6.3.268;
• Προγράμματα οδήγησης κάρτας γραφικών: ForceWare 180.60.

Για δοκιμές, ο επεξεργαστής υπερχρονίστηκε στα 3 GHz (στη διαμόρφωση 11,5x261 MHz) και στα 4 GHz (11,5x348 MHz) με συχνότητα RAM 835 MHz στην πρώτη και στη δεύτερη περίπτωση. Βίντεο σε ανάλυση Full HD 1920x1080, διάρκειας ενός λεπτού και είκοσι δευτερολέπτων. Για τη δημιουργία πρόσθετου φορτίου, ενεργοποιήθηκε ένα φίλτρο μείωσης θορύβου, οι ρυθμίσεις του οποίου παρέμειναν στις προεπιλογές.

Η κωδικοποίηση πραγματοποιήθηκε χρησιμοποιώντας τον κωδικοποιητή DivX 6.8.4. Στις ρυθμίσεις ποιότητας αυτού του κωδικοποιητή, όλες οι τιμές παραμένουν στην προεπιλογή, η λειτουργία πολλαπλών νημάτων είναι ενεργοποιημένη.

Η υποστήριξη πολλαπλών νημάτων στο TMPGEnc είναι αρχικά ενεργοποιημένη στην καρτέλα ρυθμίσεων CPU/GPU. Στην ίδια ενότητα ενεργοποιείται και το CUDA.

Όπως μπορείτε να δείτε από το παραπάνω στιγμιότυπο οθόνης, η επεξεργασία φίλτρου με χρήση CUDA είναι ενεργοποιημένη, αλλά ο αποκωδικοποιητής βίντεο υλικού δεν είναι ενεργοποιημένος. Η τεκμηρίωση του προγράμματος προειδοποιεί ότι η ενεργοποίηση της τελευταίας παραμέτρου αυξάνει τον χρόνο επεξεργασίας του αρχείου.

Με βάση τα αποτελέσματα των δοκιμών, προέκυψαν τα ακόλουθα δεδομένα:

Στα 4 GHz με ενεργοποιημένο το CUDA, κερδίσαμε μόνο μερικά δευτερόλεπτα (ή 2%), κάτι που δεν είναι ιδιαίτερα εντυπωσιακό. Αλλά σε χαμηλότερη συχνότητα, η αύξηση από την ενεργοποίηση αυτής της τεχνολογίας σάς επιτρέπει να εξοικονομήσετε περίπου 13% του χρόνου, κάτι που θα είναι αρκετά αισθητό κατά την επεξεργασία μεγάλων αρχείων. Αλλά και πάλι τα αποτελέσματα δεν είναι τόσο εντυπωσιακά όσο αναμενόταν.

Το πρόγραμμα TMPGEnc έχει δείκτη φόρτωσης CPU και CUDA σε αυτήν τη δοκιμαστική διαμόρφωση, έδειξε το φορτίο της CPU στο 20% περίπου και τον πυρήνα γραφικών στο υπόλοιπο 80%. Ως αποτέλεσμα, έχουμε το ίδιο 100% όπως όταν κάνουμε μετατροπή χωρίς CUDA και μπορεί να μην υπάρχει καθόλου χρονική διαφορά (αλλά εξακολουθεί να υπάρχει). Η μικρή χωρητικότητα μνήμης των 256 MB δεν αποτελεί επίσης περιοριστικό παράγοντα. Κρίνοντας από τις μετρήσεις από το RivaTuner, δεν χρησιμοποιήθηκαν περισσότερα από 154 MB μνήμης βίντεο κατά τη λειτουργία.

συμπεράσματα

Το πρόγραμμα TMPGEnc είναι ένα από αυτά που εισάγουν την τεχνολογία CUDA στις μάζες. Η χρήση της GPU σε αυτό το πρόγραμμα σάς επιτρέπει να επιταχύνετε τη διαδικασία επεξεργασίας βίντεο και να ανακουφίσετε σημαντικά τον κεντρικό επεξεργαστή, γεγονός που θα επιτρέψει στον χρήστη να κάνει άνετα άλλες εργασίες ταυτόχρονα. Στο συγκεκριμένο παράδειγμά μας, η κάρτα γραφικών GeForce 8800GT 256MB βελτίωσε ελαφρώς την απόδοση χρονισμού κατά τη μετατροπή βίντεο με βάση έναν υπερχρονισμένο επεξεργαστή Pentium Dual-Core E5200. Αλλά είναι ξεκάθαρα ορατό ότι όσο μειώνεται η συχνότητα, αυξάνεται το κέρδος από την ενεργοποίηση του CUDA σε ασθενείς επεξεργαστές, το κέρδος από τη χρήση του θα είναι πολύ μεγαλύτερο. Στο πλαίσιο αυτής της εξάρτησης, είναι πολύ λογικό να υποθέσουμε ότι ακόμη και με αύξηση του φορτίου (για παράδειγμα, η χρήση ενός πολύ μεγάλου αριθμού πρόσθετων φίλτρων βίντεο), τα αποτελέσματα ενός συστήματος με CUDA θα διακρίνονται από σημαντικό δέλτα της διαφοράς στο χρόνο που αφιερώθηκε στη διαδικασία κωδικοποίησης. Επίσης, μην ξεχνάτε ότι το G92 δεν είναι το πιο ισχυρό τσιπ αυτή τη στιγμή και οι πιο σύγχρονες κάρτες γραφικών θα παρέχουν σημαντικά υψηλότερη απόδοση σε τέτοιες εφαρμογές. Ωστόσο, ενώ η εφαρμογή εκτελείται, η GPU δεν είναι πλήρως φορτωμένη και, πιθανώς, η κατανομή του φορτίου εξαρτάται από κάθε διαμόρφωση ξεχωριστά, συγκεκριμένα από τον συνδυασμό επεξεργαστή/κάρτας βίντεο, ο οποίος τελικά μπορεί να δώσει μεγαλύτερη (ή μικρότερη) αύξηση ως ποσοστό ενεργοποίησης CUDA. Σε κάθε περίπτωση, για όσους εργάζονται με μεγάλους όγκους δεδομένων βίντεο, αυτή η τεχνολογία θα τους επιτρέψει να εξοικονομήσουν σημαντικά χρόνο.

Είναι αλήθεια ότι το CUDA δεν έχει ακόμη αποκτήσει ευρεία δημοτικότητα η ποιότητα του λογισμικού που λειτουργεί με αυτήν την τεχνολογία απαιτεί βελτίωση. Στο πρόγραμμα TMPGEnc 4.0 XPress που εξετάσαμε, αυτή η τεχνολογία δεν λειτουργούσε πάντα. Το ίδιο βίντεο θα μπορούσε να επανακωδικοποιηθεί πολλές φορές και, στη συνέχεια, ξαφνικά, την επόμενη φορά που κυκλοφόρησε, το φορτίο CUDA ήταν ήδη 0%. Και αυτό το φαινόμενο ήταν εντελώς τυχαίο σε εντελώς διαφορετικά λειτουργικά συστήματα. Επίσης, το εν λόγω πρόγραμμα αρνήθηκε να χρησιμοποιήσει το CUDA κατά την κωδικοποίηση σε μορφή XviD, αλλά δεν υπήρχαν προβλήματα με τον δημοφιλή κωδικοποιητή DivX.

Ως αποτέλεσμα, μέχρι στιγμής η τεχνολογία CUDA μπορεί να αυξήσει σημαντικά την απόδοση των προσωπικών υπολογιστών μόνο σε ορισμένες εργασίες. Αλλά το πεδίο εφαρμογής μιας τέτοιας τεχνολογίας θα επεκταθεί και η διαδικασία αύξησης του αριθμού των πυρήνων σε συμβατικούς επεξεργαστές υποδηλώνει αύξηση της ζήτησης για παράλληλους υπολογιστές πολλαπλών νημάτων σε σύγχρονες εφαρμογές λογισμικού. Δεν είναι τυχαίο που πρόσφατα όλοι οι ηγέτες του κλάδου έχουν εμμονή με την ιδέα του συνδυασμού CPU και GPU σε μια ενοποιημένη αρχιτεκτονική (απλώς θυμηθείτε το πολυδιαφημισμένο AMD Fusion). Ίσως το CUDA να είναι ένα από τα στάδια της διαδικασίας αυτής της ενοποίησης.

Ευχαριστούμε τις ακόλουθες εταιρείες για την παροχή εξοπλισμού δοκιμής:

Σύμφωνα με τη θεωρία της εξέλιξης του Δαρβίνου, ο πρώτος πίθηκος (αν
για την ακρίβεια - homo antecessor, human predecessor) αργότερα μετατράπηκε σε
σε μας. Κέντρα υπολογιστών πολλών τόνων με χίλιους ή περισσότερους ραδιοσωλήνες,
που καταλάμβαναν ολόκληρα δωμάτια αντικαταστάθηκαν από φορητούς υπολογιστές μισού κιλού, οι οποίοι παρεμπιπτόντως,
δεν θα είναι κατώτερο σε απόδοση από το πρώτο. Οι γραφομηχανές κατά του κατακλυσμού έχουν γίνει
στην εκτύπωση οτιδήποτε και σε οτιδήποτε (ακόμα και στο ανθρώπινο σώμα)
πολυλειτουργικές συσκευές. Οι γίγαντες των επεξεργαστών αποφάσισαν ξαφνικά να τείχη
πυρήνας γραφικών σε "πέτρα". Και οι κάρτες γραφικών άρχισαν όχι μόνο να δείχνουν μια εικόνα με
αποδεκτή ποιότητα FPS και γραφικών, αλλά και εκτέλεση κάθε είδους υπολογισμών. Ναί
ακόμα πώς να παράγει! Θα συζητηθεί η τεχνολογία των υπολογιστών πολλαπλών νημάτων με χρήση GPU.

Γιατί GPU;

Αναρωτιέμαι γιατί αποφάσισαν να μεταφέρουν όλη την υπολογιστική ισχύ στα γραφικά
προσαρμογέας? Όπως μπορείτε να δείτε, οι επεξεργαστές εξακολουθούν να είναι στη μόδα και είναι απίθανο να εγκαταλείψουν τη ζεστασιά τους
θέση. Αλλά η GPU έχει μερικά ατού στο μανίκι της, μαζί με ένα τζόκερ και πολλά μανίκια
αρκετά. Ένας σύγχρονος κεντρικός επεξεργαστής έχει σχεδιαστεί για να επιτυγχάνει το μέγιστο
απόδοση κατά την επεξεργασία δεδομένων ακεραίων και κινητής υποδιαστολής
κόμμα, χωρίς να ανησυχούμε ιδιαίτερα για την παράλληλη επεξεργασία πληροφοριών. Στο ίδιο
χρόνο, η αρχιτεκτονική της κάρτας βίντεο σάς επιτρέπει να "παραλληλίσετε" γρήγορα και χωρίς προβλήματα
επεξεργασία δεδομένων. Από τη μία, υπολογίζονται τα πολύγωνα (λόγω του τρισδιάστατου μεταφορέα),
από την άλλη, επεξεργασία υφής pixel. Είναι σαφές ότι υπάρχει μια «αρμονική
κατανομή» του φορτίου στον πυρήνα της κάρτας. Επιπλέον, απόδοση μνήμης και επεξεργαστή βίντεο
πιο βέλτιστο από τον συνδυασμό "RAM-cache-processor". Η στιγμή μιας μονάδας δεδομένων
στην κάρτα γραφικών αρχίζει να υποβάλλεται σε επεξεργασία από έναν επεξεργαστή ροής GPU, έναν άλλο
Η μονάδα φορτώνεται παράλληλα σε μια άλλη και, καταρχήν, είναι εύκολο να επιτευχθεί
Φορτίο GPU συγκρίσιμο με το εύρος ζώνης διαύλου,
Ωστόσο, για να συμβεί αυτό, οι μεταφορείς πρέπει να φορτώνονται ομοιόμορφα, χωρίς
τυχόν μεταβάσεις και διακλαδώσεις υπό όρους. Ο κεντρικός επεξεργαστής, λόγω του
Η ευελιξία απαιτεί μια πλήρη κρυφή μνήμη για τις ανάγκες επεξεργασίας της
πληροφορίες.

Οι ειδικοί έχουν σκεφτεί τη δουλειά των GPU στον παράλληλο υπολογισμό και
μαθηματικά και κατέληξε στη θεωρία ότι πολλοί επιστημονικοί υπολογισμοί είναι από πολλές απόψεις παρόμοιοι με
Επεξεργασία τρισδιάστατων γραφικών. Πολλοί ειδικοί πιστεύουν ότι ο θεμελιώδης παράγοντας στην
ανάπτυξη GPGPU (Υπολογισμός γενικού σκοπού σε GPU – καθολική
υπολογισμοί με χρήση κάρτας βίντεο) ήταν η εμφάνιση του έργου Brook GPU το 2003.

Οι δημιουργοί του έργου από το Πανεπιστήμιο του Στάνφορντ έπρεπε να λύσουν ένα δύσκολο
πρόβλημα: υλικό και λογισμικό για να αναγκάσουν τον προσαρμογέα γραφικών να παράγει
διαφορετικούς υπολογισμούς. Και τα κατάφεραν. Χρησιμοποιώντας τη γενική γλώσσα C,
Αμερικανοί επιστήμονες έκαναν τη GPU να λειτουργεί σαν επεξεργαστής, προσαρμοσμένη
παράλληλη επεξεργασία. Μετά το Brook, εμφανίστηκαν μια σειρά από έργα σε υπολογισμούς VGA,
όπως Accelerator library, Brahma library, system
Μεταπρογραμματισμός GPU++ και άλλα.

CUDA!

Το προαίσθημα των προοπτικών ανάπτυξης αναγκαστική AMDΚαι NVIDIA
κολλήστε στην GPU του Brook σαν πίτμπουλ. Αν παραλείψουμε την πολιτική μάρκετινγκ, τότε
Εφαρμόζοντας τα πάντα σωστά, μπορείτε να αποκτήσετε βάση όχι μόνο στον τομέα των γραφικών
αγορά, αλλά και στους υπολογιστές (δείτε ειδικές κάρτες υπολογιστών και
διακομιστές Teslaμε εκατοντάδες πολυεπεξεργαστές), εκτοπίζοντας τις συνηθισμένες CPU.

Όπως ήταν φυσικό, οι «άρχοντες του FPS» χώρισαν τους δρόμους τους στο εμπόδιο, ο καθένας με τον δικό του τρόπο.
διαδρομή, αλλά η βασική αρχή παρέμεινε αμετάβλητη - να κάνετε υπολογισμούς
χρησιμοποιώντας GPU. Και τώρα θα ρίξουμε μια πιο προσεκτική ματιά στην «πράσινη» τεχνολογία - CUDA
(Υπολογισμός Ενοποιημένης Αρχιτεκτονικής Συσκευών).

Η δουλειά της «ηρωίδας» μας είναι να παρέχει ένα API, δύο ταυτόχρονα.
Το πρώτο είναι υψηλού επιπέδου, το CUDA Runtime, το οποίο αντιπροσωπεύει τις λειτουργίες που
αναλύονται σε απλούστερα επίπεδα και περνούν στο κατώτερο API - CUDA Driver. Έτσι
ότι η φράση "υψηλού επιπέδου" είναι μια έκταση που πρέπει να εφαρμοστεί στη διαδικασία. Όλο το αλάτι είναι
ακριβώς στο πρόγραμμα οδήγησης και οι βιβλιοθήκες που δημιουργήθηκαν ευγενικά θα σας βοηθήσουν να το αποκτήσετε
προγραμματιστές NVIDIA: CUBLAS (εργαλεία για μαθηματικούς υπολογισμούς) και
FFT (υπολογισμός με χρήση του αλγόριθμου Fourier). Λοιπόν, ας περάσουμε στο πρακτικό
μέρη του υλικού.

Ορολογία CUDA

NVIDIAλειτουργεί με πολύ μοναδικούς ορισμούς για το CUDA API. Αυτοί
διαφέρουν από τους ορισμούς που χρησιμοποιούνται για την εργασία με έναν κεντρικό επεξεργαστή.

Νήμα– ένα σύνολο δεδομένων που πρέπει να υποβληθούν σε επεξεργασία (όχι
απαιτεί μεγάλους πόρους επεξεργασίας).

Στημόνι– μια ομάδα 32 νημάτων. Τα δεδομένα υποβάλλονται σε επεξεργασία μόνο
παραμορφώνει, επομένως ένα στραβό είναι ο ελάχιστος όγκος δεδομένων.

ΟΙΚΟΔΟΜΙΚΟ ΤΕΤΡΑΓΩΝΟ– ένα σύνολο ροών (από 64 έως 512) ή ένα σύνολο
στημόνι (από 2 έως 16).

Πλέγμαείναι μια συλλογή από μπλοκ. Αυτή η διαίρεση δεδομένων
χρησιμοποιείται αποκλειστικά για τη βελτίωση της απόδοσης. Έτσι, εάν ο αριθμός
οι πολυεπεξεργαστές είναι μεγάλοι, τότε τα μπλοκ θα εκτελεστούν παράλληλα. Αν με
καμία τύχη με την κάρτα (οι προγραμματιστές συνιστούν τη χρήση
προσαρμογέας όχι χαμηλότερος από GeForce 8800 GTS 320 MB), τότε τα μπλοκ δεδομένων θα υποβληθούν σε επεξεργασία
διαδοχικά.

Η NVIDIA εισάγει επίσης έννοιες όπως πυρήνας, πλήθος
Και συσκευή.

Δουλεύουμε!

Για να εργαστείτε πλήρως με το CUDA χρειάζεστε:

1. Γνωρίστε τη δομή των πυρήνων shader της GPU, από την ουσία του προγραμματισμού
συνίσταται στην ομοιόμορφη κατανομή του φορτίου μεταξύ τους.
2. Να είναι σε θέση να προγραμματίζει στο περιβάλλον C, λαμβάνοντας υπόψη ορισμένες πτυχές.

προγραμματιστές NVIDIAαποκάλυψε τα «μέσα» της κάρτας βίντεο αρκετές φορές
διαφορετικά από ό,τι έχουμε συνηθίσει να βλέπουμε. Οπότε, θέλοντας και μη, θα πρέπει να μελετήσετε τα πάντα
λεπτότητες της αρχιτεκτονικής. Ας δούμε τη δομή της θρυλικής "πέτρας" G80 GeForce 8800
GTX.

Ο πυρήνας shader αποτελείται από οκτώ συμπλέγματα TPC (Texture Processor Cluster)
επεξεργαστές υφής (έτσι, GeForce GTX 280– 15 πυρήνες, 8800 GTS
είναι έξι από αυτούς 8600 – τέσσερα, κλπ.). Αυτά με τη σειρά τους αποτελούνται από δύο
πολυεπεξεργαστές ροής (εφεξής SM). SM (όλοι τους
16) αποτελείται από το μπροστινό μέρος (λύνει τα προβλήματα ανάγνωσης και αποκωδικοποίησης οδηγιών) και
σωληνώσεις πίσω άκρου (τελική έξοδος οδηγιών), καθώς και οκτώ βαθμωτές SP (shader
επεξεργαστή) και δύο SFU (μονάδες υπερλειτουργίας). Για κάθε χτύπημα (μονάδα
time) το μπροστινό μέρος επιλέγει το στημόνι και το επεξεργάζεται. Έτσι ώστε όλο το στημόνι να ρέει
(να θυμίσω, είναι 32) επεξεργασμένα, απαιτούνται 32/8 = 4 κύκλοι στο τέλος του μεταφορέα.

Κάθε πολυεπεξεργαστής έχει αυτό που ονομάζεται κοινόχρηστη μνήμη.
Το μέγεθός του είναι 16 kilobyte και παρέχει στον προγραμματιστή πλήρη ελευθερία
Ενέργειες. Διαδώστε όπως θέλετε :). Η κοινή μνήμη παρέχει επικοινωνία μεταξύ των νημάτων
ένα μπλοκ και δεν προορίζεται να λειτουργήσει με σκίαστρους pixel.

Τα SM μπορούν επίσης να έχουν πρόσβαση στο GDDR. Για να γίνει αυτό, τους δόθηκαν 8 kilobyte το καθένα.
μνήμη cache που αποθηκεύει όλα τα πιο σημαντικά πράγματα για εργασία (για παράδειγμα, υπολογιστές
σταθερές).

Ο πολυεπεξεργαστής έχει 8192 καταχωρητές. Ο αριθμός των ενεργών μπλοκ δεν μπορεί να είναι
περισσότερα από οκτώ και ο αριθμός των στημονιών δεν είναι μεγαλύτερος από 768/32 = 24. Από αυτό είναι σαφές ότι το G80
μπορεί να επεξεργαστεί το πολύ 32*16*24 = 12288 νήματα ανά μονάδα χρόνου. Δεν μπορείς παρά
λάβετε υπόψη αυτούς τους αριθμούς κατά τη βελτιστοποίηση του προγράμματος στο μέλλον (σε μία κλίμακα
– μέγεθος μπλοκ, από την άλλη – ο αριθμός των νημάτων). Η ισορροπία παραμέτρων μπορεί να παίξει κάποιο ρόλο
σημαντικό ρόλο λοιπόν στο μέλλον NVIDIAσυνιστά τη χρήση μπλοκ
με 128 ή 256 νήματα. Ένα μπλοκ 512 νημάτων είναι αναποτελεσματικό επειδή έχει
αυξημένες καθυστερήσεις. Λαμβάνοντας υπόψη όλες τις λεπτές αποχρώσεις της δομής της κάρτας βίντεο GPU συν
καλές δεξιότητες προγραμματισμού, μπορείτε να δημιουργήσετε πολύ παραγωγικές
εργαλείο για παράλληλους υπολογιστές. Παρεμπιπτόντως, για τον προγραμματισμό...

Προγραμματισμός

Για «δημιουργικότητα» με το CUDA χρειάζεστε Η κάρτα γραφικών GeForce δεν είναι χαμηλότερη
επεισόδιο όγδοο. ΜΕ

Επίσημος ιστότοπος πρέπει να κατεβάσετε τρία πακέτα λογισμικού: πρόγραμμα οδήγησης από
Υποστήριξη CUDA (κάθε λειτουργικό σύστημα έχει το δικό του), το ίδιο το πακέτο CUDA SDK (το δεύτερο
έκδοση beta) και πρόσθετες βιβλιοθήκες (εργαλειοθήκη CUDA). Η τεχνολογία υποστηρίζει
λειτουργικά συστήματα Windows (XP και Vista), Linux και Mac OS X. Να μελετήσω Ι
επέλεξε Vista Ultimate Edition x64 (κοιτάζοντας μπροστά, θα πω ότι το σύστημα συμπεριφέρθηκε
Τέλεια). Την εποχή που γράφονταν αυτές οι γραμμές, ήταν σχετικό με τη δουλειά
ForceWare πρόγραμμα οδήγησης 177.35. Χρησιμοποιείται ως σύνολο εργαλείων
Πακέτο λογισμικού Borland C++ 6 Builder (αν και οποιοδήποτε περιβάλλον λειτουργεί με
γλώσσα Γ).

Θα είναι εύκολο για ένα άτομο που γνωρίζει τη γλώσσα να συνηθίσει σε ένα νέο περιβάλλον. Το μόνο που απαιτείται είναι
θυμηθείτε τις βασικές παραμέτρους. Λέξη-κλειδί _global_ (τοποθετείται πριν από τη συνάρτηση)
υποδεικνύει ότι η συνάρτηση ανήκει στον πυρήνα. Θα την καλέσει το κεντρικό
επεξεργαστή και όλη η δουλειά θα γίνει στη GPU. Η κλήση _global_ απαιτεί περισσότερα
συγκεκριμένες λεπτομέρειες, δηλαδή μέγεθος πλέγματος, μέγεθος μπλοκ και ποιος θα είναι ο πυρήνας
εφαρμοσμένος. Για παράδειγμα, η γραμμή _global_ void saxpy_parallel<<>>, όπου X –
το μέγεθος του πλέγματος, και το Y είναι το μέγεθος του μπλοκ, καθορίζει αυτές τις παραμέτρους.

Το σύμβολο _device_ σημαίνει ότι η συνάρτηση θα κληθεί από τον πυρήνα γραφικών, γνωστό και ως
θα ακολουθήσει όλες τις οδηγίες. Αυτή η λειτουργία βρίσκεται στη μνήμη του πολυεπεξεργαστή,
Ως εκ τούτου, είναι αδύνατο να λάβετε τη διεύθυνσή της. Το πρόθεμα _host_ σημαίνει ότι η κλήση
και η επεξεργασία θα γίνεται μόνο με τη συμμετοχή της CPU. Πρέπει να ληφθεί υπόψη ότι _παγκόσμια_ και
Οι _συσκευές_ δεν μπορούν να καλούν ο ένας τον άλλον και δεν μπορούν να καλούν τον εαυτό τους.

Επίσης, η γλώσσα για το CUDA έχει μια σειρά από λειτουργίες για εργασία με μνήμη βίντεο: cudafree
(απελευθέρωση μνήμης μεταξύ GDDR και RAM), cudamemcpy και cudamemcpy2D (αντιγραφή
μνήμη μεταξύ GDDR και RAM) και cudamalloc (εκχώρηση μνήμης).

Όλοι οι κωδικοί προγράμματος μεταγλωττίζονται από το CUDA API. Πρώτα λαμβάνεται
κωδικός που προορίζεται αποκλειστικά για τον κεντρικό επεξεργαστή και υπόκειται σε
τυπική μεταγλώττιση και άλλος κώδικας που προορίζεται για τον προσαρμογέα γραφικών
ξαναγραφτεί στην ενδιάμεση γλώσσα PTX (σαν assembler) για
εντοπισμός πιθανών σφαλμάτων. Μετά από όλους αυτούς τους «χορούς» ο τελικός
μετάφραση (μετάφραση) εντολών σε γλώσσα κατανοητή για GPU/CPU.

Κιτ μελέτης

Σχεδόν όλες οι πτυχές του προγραμματισμού περιγράφονται στην τεκμηρίωση που ακολουθεί
μαζί με το πρόγραμμα οδήγησης και δύο εφαρμογές, καθώς και στον ιστότοπο των προγραμματιστών. Μέγεθος
το άρθρο δεν αρκεί για να τα περιγράψει (ο ενδιαφερόμενος αναγνώστης θα πρέπει να επισυνάψει
λίγη προσπάθεια και μελετήστε μόνοι σας το υλικό).

Το πρόγραμμα περιήγησης CUDA SDK αναπτύχθηκε ειδικά για αρχάριους. Οποιοσδήποτε μπορεί
αισθανθείτε τη δύναμη του παράλληλου υπολογισμού από πρώτο χέρι (το καλύτερο τεστ για
σταθερότητα – τα παραδείγματα λειτουργούν χωρίς τεχνουργήματα ή συντριβές). Η εφαρμογή έχει
μεγάλος αριθμός ενδεικτικών μίνι προγραμμάτων (61 «τεστ»). Για κάθε εμπειρία υπάρχει
Αναλυτική τεκμηρίωση του κώδικα του προγράμματος συν αρχεία PDF. Είναι αμέσως φανερό ότι οι άνθρωποι
όσοι είναι παρόντες με τις δημιουργίες τους στο πρόγραμμα περιήγησης κάνουν σοβαρή δουλειά.
Μπορείτε επίσης να συγκρίνετε την ταχύτητα του επεξεργαστή και της κάρτας βίντεο κατά την επεξεργασία
δεδομένα. Για παράδειγμα, σάρωση πολυδιάστατων συστοιχιών με κάρτα βίντεο GeForce 8800
GTΠαράγει 512 MB με ένα μπλοκ με 256 νήματα σε 0,17109 χιλιοστά του δευτερολέπτου.
Η τεχνολογία δεν αναγνωρίζει SLI tandems, οπότε αν έχετε ένα δίδυμο ή τρίο,
απενεργοποιήστε τη λειτουργία "σύζευξης" πριν εργαστείτε, διαφορετικά το CUDA θα δει μόνο μία
συσκευή Διπύρηνο AMD Athlon 64 X2(συχνότητα πυρήνα 3000 MHz) ίδια εμπειρία
περνά σε 2,761528 χιλιοστά του δευτερολέπτου. Αποδεικνύεται ότι το G92 είναι περισσότερο από 16 φορές
πιο γρήγορα από έναν βράχο AMD! Όπως μπορείτε να δείτε, αυτό απέχει πολύ από ένα ακραίο σύστημα
σε συνδυασμό με ένα λειτουργικό σύστημα που δεν αγαπήθηκε από τις μάζες δείχνει καλό
Αποτελέσματα.

Εκτός από το πρόγραμμα περιήγησης, υπάρχει μια σειρά από προγράμματα χρήσιμα για την κοινωνία. Πλίθα
προσάρμοσε τα προϊόντα της στη νέα τεχνολογία. Τώρα το Photoshop CS4 είναι πλήρες
χρησιμοποιεί τουλάχιστον τους πόρους των προσαρμογέων γραφικών (πρέπει να κατεβάσετε ένα ειδικό
συνδέω). Με προγράμματα όπως ο μετατροπέας πολυμέσων Badaboom και το RapiHD μπορείτε
αποκωδικοποίηση βίντεο σε μορφή MPEG-2. Καλό για επεξεργασία ήχου
Το δωρεάν βοηθητικό πρόγραμμα Accelero είναι κατάλληλο. Η ποσότητα του λογισμικού που είναι προσαρμοσμένη για το CUDA API,
αναμφίβολα θα αυξηθεί.

Και αυτή την ώρα...

Εν τω μεταξύ, διαβάζετε αυτό το υλικό, σκληρά εργαζόμενοι από τις ανησυχίες του επεξεργαστή
αναπτύσσουν τις δικές τους τεχνολογίες για την ενσωμάτωση των GPU σε CPU. Απο έξω AMDΟλα
είναι σαφές: έχουν τεράστια εμπειρία που αποκτήθηκε μαζί με ATI.

Η δημιουργία «μικροσυσκευών», Fusion, θα αποτελείται από αρκετούς πυρήνες κάτω
κωδική ονομασία Μπουλντόζα και τσιπ βίντεο RV710 (Kong). Η σχέση τους θα είναι
πραγματοποιείται μέσω του βελτιωμένου λεωφορείου HyperTransport. Εξαρτάται από
αριθμός πυρήνων και τα χαρακτηριστικά συχνότητάς τους Η AMD σχεδιάζει να δημιουργήσει μια ολόκληρη τιμή
ιεραρχία των «λίθων». Προβλέπεται επίσης η παραγωγή επεξεργαστών για φορητούς υπολογιστές (Falcon),
και για gadget πολυμέσων (Bobcat). Επιπλέον, είναι η εφαρμογή της τεχνολογίας
στις φορητές συσκευές θα είναι η αρχική πρόκληση για τους Καναδούς. Με ανάπτυξη
παράλληλοι υπολογιστές, η χρήση τέτοιων «λίθων» θα πρέπει να είναι πολύ δημοφιλής.

Intelλίγο πίσω στον χρόνο με τον Larrabee του. Προϊόντα AMD,
αν δεν συμβεί τίποτα, θα εμφανιστούν στα ράφια των καταστημάτων στα τέλη του 2009 - αρχές
2010. Και η απόφαση του εχθρού θα έρθει στο φως μόνο σε σχεδόν δύο
της χρονιάς.

Ο Larrabee θα έχει μεγάλο αριθμό (διαβάστε: εκατοντάδες) πυρήνες. Στην αρχή
Θα υπάρχουν επίσης προϊόντα σχεδιασμένα για 8 – 64 πυρήνες. Μοιάζουν πολύ με το Pentium, αλλά
αρκετά βαριά ανακατασκευή. Κάθε πυρήνας έχει 256 kilobyte μνήμης cache L2
(το μέγεθός του θα αυξηθεί με την πάροδο του χρόνου). Η σχέση θα πραγματοποιηθεί μέσω
Αμφίδρομος δακτύλιος 1024-bit. Η Intel λέει ότι το "παιδί" της θα είναι
λειτουργούν τέλεια με το DirectX και το Open GL API (για προγραμματιστές της Apple), οπότε όχι
δεν απαιτείται παρέμβαση λογισμικού.

Γιατί σας τα είπα όλα αυτά; Είναι προφανές ότι οι Larrabee και Fusion δεν θα εκτοπίσουν
κανονικούς, σταθερούς επεξεργαστές από την αγορά, όπως δεν θα εξαναγκαστούν να βγουν από την αγορά
κάρτες γραφικών. Για τους παίκτες και τους λάτρεις των extreme sports, το απόλυτο όνειρο θα παραμείνει
CPU πολλαπλών πυρήνων και μια σειρά από πολλά κορυφαία VGA. Αλλά τι ακόμη
Οι εταιρείες επεξεργασίας στρέφονται στον παράλληλο υπολογισμό με βάση τις αρχές
παρόμοιο με το GPGPU, λέει πολλά. Ειδικότερα, για το τι τέτοια
τεχνολογία όπως το CUDA έχει δικαίωμα ύπαρξης και, πιθανότατα, θα υπάρχει
πολύ δημοφιλής.

Μια σύντομη περίληψη

Ο παράλληλος υπολογισμός με χρήση κάρτας βίντεο είναι απλώς ένα καλό εργαλείο
στα χέρια ενός εργατικού προγραμματιστή. Σχεδόν οι μεταποιητές καθοδηγούνται από το νόμο του Μουρ
θα έρθει το τέλος. Εταιρείες NVIDIAυπάρχει πολύς δρόμος ακόμα
προωθώντας το API του στις μάζες (το ίδιο μπορεί να ειπωθεί για το πνευματικό τέκνο ATI/AMD).
Το πώς θα είναι θα δείξει το μέλλον. Οπότε το CUDA θα επιστρέψει :).

ΥΣΤΕΡΟΓΡΑΦΟ. Προτείνω να επισκεφθούν αρχάριοι προγραμματιστές και ενδιαφερόμενοι
τις ακόλουθες «εικονικές εγκαταστάσεις»:

Επίσημος ιστότοπος και ιστότοπος της NVIDIA
GPGPU.com. Ολα
οι πληροφορίες που παρέχονται είναι στα Αγγλικά, αλλά τουλάχιστον σας ευχαριστώ που δεν περιλαμβάνονται
κινέζικα Προχωρήστε λοιπόν! Ελπίζω ο συγγραφέας να σε βοήθησε έστω λίγο
συναρπαστικά ταξίδια στην εξερεύνηση CUDA!

Τεχνολογία CUDA

Βλαντιμίρ Φρόλοφ,[email προστατευμένο]

σχόλιο

Το άρθρο μιλά για την τεχνολογία CUDA, η οποία επιτρέπει σε έναν προγραμματιστή να χρησιμοποιεί κάρτες βίντεο ως ισχυρές υπολογιστικές μονάδες. Τα εργαλεία που παρέχει η Nvidia καθιστούν δυνατή τη σύνταξη προγραμμάτων μονάδας επεξεργασίας γραφικών (GPU) σε ένα υποσύνολο της γλώσσας C++. Αυτό απαλλάσσει τον προγραμματιστή από την ανάγκη χρήσης shaders και κατανόησης της λειτουργίας του αγωγού γραφικών. Το άρθρο παρέχει παραδείγματα προγραμματισμού χρησιμοποιώντας CUDA και διάφορες τεχνικές βελτιστοποίησης.

1. Εισαγωγή

Η ανάπτυξη της υπολογιστικής τεχνολογίας έχει προχωρήσει ραγδαία τις τελευταίες δεκαετίες. Τόσο γρήγορα που οι προγραμματιστές επεξεργαστών έχουν ήδη φτάσει στο λεγόμενο «αδιέξοδο πυριτίου». Η αχαλίνωτη αύξηση της συχνότητας του ρολογιού έχει καταστεί αδύνατη λόγω πολλών σοβαρών τεχνολογικών λόγων.

Αυτός είναι εν μέρει ο λόγος που όλοι οι κατασκευαστές σύγχρονων υπολογιστικών συστημάτων κινούνται προς την αύξηση του αριθμού των επεξεργαστών και των πυρήνων, αντί να αυξάνουν τη συχνότητα ενός επεξεργαστή. Ο αριθμός των πυρήνων κεντρικής μονάδας επεξεργασίας (CPU) σε προηγμένα συστήματα είναι ήδη 8.

Ένας άλλος λόγος είναι η σχετικά χαμηλή ταχύτητα της μνήμης RAM. Ανεξάρτητα από το πόσο γρήγορα λειτουργεί ο επεξεργαστής, τα σημεία συμφόρησης, όπως δείχνει η πρακτική, δεν είναι καθόλου αριθμητικές πράξεις, αλλά μάλλον ανεπιτυχείς προσβάσεις στη μνήμη - αστοχίες της προσωρινής μνήμης.

Ωστόσο, αν κοιτάξετε προς τους επεξεργαστές γραφικών GPU (Graphics Processing Unit), πήραν το δρόμο του παραλληλισμού πολύ νωρίτερα. Στις σημερινές κάρτες γραφικών, για παράδειγμα στην GF8800GTX, ο αριθμός των επεξεργαστών μπορεί να φτάσει τους 128. Η απόδοση τέτοιων συστημάτων, με επιδέξιο προγραμματισμό, μπορεί να είναι αρκετά σημαντική (Εικ. 1).

Ρύζι. 1. Αριθμός λειτουργιών κινητής υποδιαστολής για CPU και GPU

Όταν εμφανίστηκαν για πρώτη φορά στην αγορά οι πρώτες κάρτες γραφικών, ήταν αρκετά απλές (σε σύγκριση με τον κεντρικό επεξεργαστή) εξαιρετικά εξειδικευμένες συσκευές σχεδιασμένες να απαλλάσσουν τον επεξεργαστή από το βάρος της οπτικοποίησης δισδιάστατων δεδομένων. Με την ανάπτυξη της βιομηχανίας τυχερών παιχνιδιών και την εμφάνιση τέτοιων τρισδιάστατων παιχνιδιών όπως το Doom (Εικ. 2) και το Wolfenstein 3D (Εικ. 3), προέκυψε η ανάγκη για τρισδιάστατη απεικόνιση.

Εικόνες 2,3. Παιχνίδια Doom και Wolfenstein 3D

Από τη στιγμή που η 3Dfx δημιούργησε τις πρώτες κάρτες βίντεο Voodoo (1996) έως το 2001, μόνο ένα σταθερό σύνολο λειτουργιών σε δεδομένα εισόδου εφαρμόστηκε στη GPU.

Οι προγραμματιστές δεν είχαν άλλη επιλογή στον αλγόριθμο απόδοσης και για να αυξήσουν την ευελιξία, εμφανίστηκαν shaders - μικρά προγράμματα που εκτελούνται από την κάρτα βίντεο για κάθε κορυφή ή για κάθε pixel. Οι εργασίες τους περιελάμβαναν μετασχηματισμούς σε κορυφές και σκίαση—υπολογίζοντας τον φωτισμό σε ένα σημείο, για παράδειγμα, χρησιμοποιώντας το μοντέλο Phong.

Αν και τα shader έχουν προχωρήσει πολύ αυτές τις μέρες, θα πρέπει να γίνει κατανοητό ότι σχεδιάστηκαν για εξαιρετικά εξειδικευμένες εργασίες όπως ο τρισδιάστατος μετασχηματισμός και η ραστεροποίηση. Ενώ οι GPU εξελίσσονται προς συστήματα πολλαπλών επεξεργαστών γενικής χρήσης, οι γλώσσες shader παραμένουν εξαιρετικά εξειδικευμένες.

Μπορούν να συγκριθούν με το FORTRAN με την έννοια ότι, όπως και το FORTRAN, ήταν τα πρώτα, αλλά σχεδιάστηκαν για να λύσουν μόνο έναν τύπο προβλήματος. Οι σκιαδόροι χρησιμοποιούνται ελάχιστα για την επίλυση οποιωνδήποτε άλλων προβλημάτων εκτός από τους τρισδιάστατους μετασχηματισμούς και τη ραστεροποίηση, όπως το FORTRAN δεν είναι βολικό για την επίλυση προβλημάτων που δεν σχετίζονται με αριθμητικούς υπολογισμούς.

Σήμερα υπάρχει η τάση να χρησιμοποιούνται κάρτες βίντεο με ασυνήθιστο τρόπο για την επίλυση προβλημάτων στους τομείς της κβαντικής μηχανικής, της τεχνητής νοημοσύνης, των φυσικών υπολογισμών, της κρυπτογραφίας, της φυσικής σωστής απεικόνισης, της ανακατασκευής από φωτογραφίες, της αναγνώρισης κ.λπ. Αυτές οι εργασίες δεν είναι βολικό να επιλυθούν στο πλαίσιο των API γραφικών (DirectX, OpenGL), καθώς αυτά τα API δημιουργήθηκαν για εντελώς διαφορετικές εφαρμογές.

Η ανάπτυξη γενικού προγραμματισμού σε GPU (General Programming on GPU, GPGPU) οδήγησε λογικά στην εμφάνιση τεχνολογιών που στοχεύουν σε ένα ευρύτερο φάσμα εργασιών από το rasterization. Ως αποτέλεσμα, η Nvidia δημιούργησε το Compute Unified Device Architecture (ή CUDA για συντομία) και η ανταγωνιστική εταιρεία ATI δημιούργησε την τεχνολογία STREAM.

Θα πρέπει να σημειωθεί ότι τη στιγμή της συγγραφής αυτού του άρθρου, η τεχνολογία STREAM ήταν πολύ πίσω από το CUDA στην ανάπτυξη και επομένως δεν θα ληφθεί υπόψη εδώ. Θα επικεντρωθούμε στο CUDA, μια τεχνολογία GPGPU που σας επιτρέπει να γράφετε προγράμματα σε ένα υποσύνολο της γλώσσας C++.

2. Θεμελιώδης διαφορά μεταξύ CPU και GPU

Ας ρίξουμε μια γρήγορη ματιά σε μερικές από τις σημαντικές διαφορές μεταξύ των περιοχών και των δυνατοτήτων των εφαρμογών CPU και κάρτας βίντεο.

2.1. Δυνατότητες

Η CPU σχεδιάστηκε αρχικά για την επίλυση γενικών προβλημάτων και λειτουργεί με τυχαία διευθυνσιοδοτούμενη μνήμη. Τα προγράμματα στη CPU μπορούν να έχουν άμεση πρόσβαση σε οποιαδήποτε γραμμικά και ομοιογενή κελιά μνήμης.

Αυτό δεν ισχύει για τις GPU. Όπως θα μάθετε διαβάζοντας αυτό το άρθρο, το CUDA έχει έως και 6 τύπους μνήμης. Μπορείτε να διαβάσετε από οποιοδήποτε κελί που είναι φυσικά προσβάσιμο, αλλά δεν μπορείτε να γράψετε σε όλα τα κελιά. Ο λόγος είναι ότι η GPU είναι σε κάθε περίπτωση μια συγκεκριμένη συσκευή σχεδιασμένη για συγκεκριμένους σκοπούς. Αυτός ο περιορισμός εισήχθη για να αυξήσει την ταχύτητα ορισμένων αλγορίθμων και να μειώσει το κόστος του εξοπλισμού.

2.2. Απόδοση μνήμης

Ένα διαχρονικό πρόβλημα με τα περισσότερα υπολογιστικά συστήματα είναι ότι η μνήμη είναι πιο αργή από τον επεξεργαστή. Οι κατασκευαστές CPU λύνουν αυτό το πρόβλημα εισάγοντας κρυφές μνήμες. Οι πιο συχνά χρησιμοποιούμενες περιοχές της μνήμης τοποθετούνται στη μνήμη ή στη μνήμη cache, λειτουργώντας στη συχνότητα του επεξεργαστή. Αυτό σας επιτρέπει να εξοικονομήσετε χρόνο κατά την πρόσβαση στα δεδομένα που χρησιμοποιούνται πιο συχνά και να φορτώσετε τον επεξεργαστή με τους πραγματικούς υπολογισμούς.

Σημειώστε ότι οι κρυφές μνήμες είναι ουσιαστικά διαφανείς στον προγραμματιστή. Τόσο κατά την ανάγνωση όσο και κατά την εγγραφή, τα δεδομένα δεν πηγαίνουν απευθείας στη μνήμη RAM, αλλά περνούν από τις κρυφές μνήμες. Αυτό επιτρέπει, συγκεκριμένα, τη γρήγορη ανάγνωση μιας τιμής αμέσως μετά τη σύνταξή της.

Οι GPU (εδώ εννοούμε τις κάρτες γραφικών όγδοης σειράς GF) έχουν επίσης κρυφές μνήμες, και είναι επίσης σημαντικές, αλλά αυτός ο μηχανισμός δεν είναι τόσο ισχυρός όσο σε μια CPU. Πρώτον, δεν αποθηκεύονται όλοι οι τύποι μνήμης στην κρυφή μνήμη και, δεύτερον, οι κρυφές μνήμες είναι μόνο για ανάγνωση.

Στις GPU, οι αργές προσβάσεις στη μνήμη αποκρύπτονται χρησιμοποιώντας παράλληλους υπολογιστές. Ενώ ορισμένες εργασίες περιμένουν δεδομένα, άλλες λειτουργούν, έτοιμες για υπολογισμούς. Αυτή είναι μια από τις βασικές αρχές του CUDA, που μπορεί να βελτιώσει σημαντικά την απόδοση του συστήματος στο σύνολό του.

3. CUDA Core

3.1. Μοντέλο ροής

Η αρχιτεκτονική υπολογιστών CUDA βασίζεται στην ιδέαμία εντολή για πολλά δεδομένα(Single Instruction Multiple Data, SIMD) και έννοια πολυεπεξεργαστής.

Η έννοια SIMD σημαίνει ότι μια εντολή μπορεί να επεξεργαστεί πολλά δεδομένα ταυτόχρονα. Για παράδειγμα, η εντολή addps στους επεξεργαστές Pentium 3 και στους νεότερους επεξεργαστές Pentium σάς επιτρέπει να προσθέσετε ταυτόχρονα 4 αριθμούς κινητής υποδιαστολής απλής ακρίβειας.

Ένας πολυεπεξεργαστής είναι ένας πολυπύρηνος επεξεργαστής SIMD που επιτρέπει την εκτέλεση μόνο μιας εντολής σε όλους τους πυρήνες ανά πάσα στιγμή. Κάθε πυρήνας πολυεπεξεργαστή είναι βαθμωτός, δηλ. δεν υποστηρίζει καθαρές διανυσματικές πράξεις.

Πριν συνεχίσουμε, ας εισαγάγουμε μερικούς ορισμούς. Λάβετε υπόψη ότι σε αυτό το άρθρο θα εννοούμε τη συσκευή και τον κεντρικό υπολογιστή με εντελώς διαφορετικό τρόπο από αυτό που έχουν συνηθίσει οι περισσότεροι προγραμματιστές. Θα χρησιμοποιήσουμε τέτοιους όρους για να αποφύγουμε ασυμφωνίες με την τεκμηρίωση CUDA.

Στο άρθρο μας, ως συσκευή θα εννοούμε έναν προσαρμογέα βίντεο που υποστηρίζει το πρόγραμμα οδήγησης CUDA ή μια άλλη εξειδικευμένη συσκευή που έχει σχεδιαστεί για την εκτέλεση προγραμμάτων που χρησιμοποιούν CUDA (όπως η NVIDIA Tesla). Στο άρθρο μας, θα εξετάσουμε την GPU μόνο ως μια λογική συσκευή, αποφεύγοντας συγκεκριμένες λεπτομέρειες υλοποίησης.

Θα ονομάσουμε έναν κεντρικό υπολογιστή ένα πρόγραμμα στην κανονική μνήμη RAM του υπολογιστή που χρησιμοποιεί την CPU και εκτελεί λειτουργίες ελέγχου για την εργασία με τη συσκευή.

Στην πραγματικότητα, το μέρος του προγράμματός σας που τρέχει στη CPU είναιπλήθος, και η κάρτα γραφικών σας -συσκευή. Λογικά, η συσκευή μπορεί να αναπαρασταθεί ως ένα σύνολο πολυεπεξεργαστών (Εικ. 4) συν ένα πρόγραμμα οδήγησης CUDA.

Ρύζι. 4. Συσκευή

Ας υποθέσουμε ότι θέλουμε να εκτελέσουμε μια συγκεκριμένη διαδικασία στη συσκευή μας σε N νήματα (δηλαδή θέλουμε να παραλληλίσουμε το έργο της). Σύμφωνα με την τεκμηρίωση CUDA, ας ονομάσουμε αυτή τη διαδικασία πυρήνα.

Ένα χαρακτηριστικό της αρχιτεκτονικής CUDA είναι η οργάνωση μπλοκ-πλέγματος, η οποία είναι ασυνήθιστη για εφαρμογές πολλαπλών νημάτων (Εικ. 5). Σε αυτήν την περίπτωση, το πρόγραμμα οδήγησης CUDA κατανέμει ανεξάρτητα τους πόρους της συσκευής μεταξύ των νημάτων.

Ρύζι. 5. Οργάνωση νημάτων

Στο Σχ. 5. Ο πυρήνας ορίζεται ως πυρήνας. Όλα τα νήματα που εκτελούν αυτόν τον πυρήνα συνδυάζονται σε μπλοκ (Block) και τα μπλοκ, με τη σειρά τους, συνδυάζονται σε ένα πλέγμα (Grid).

Όπως φαίνεται στο Σχήμα 5, οι δισδιάστατοι δείκτες χρησιμοποιούνται για την αναγνώριση των νημάτων. Οι προγραμματιστές CUDA έχουν παράσχει τη δυνατότητα εργασίας με τρισδιάστατα, δισδιάστατα ή απλά (μονοδιάστατα) ευρετήρια, ανάλογα με το τι είναι πιο βολικό για τον προγραμματιστή.

Γενικά, οι δείκτες είναι τρισδιάστατα διανύσματα. Για κάθε νήμα, θα είναι γνωστά τα εξής: ο δείκτης νήματος μέσα στο μπλοκ threadIdx και ο δείκτης μπλοκ μέσα στο πλέγμα blockIdx. Κατά την εκκίνηση, όλα τα νήματα θα διαφέρουν μόνο σε αυτά τα ευρετήρια. Στην πραγματικότητα, μέσω αυτών των ευρετηρίων ο προγραμματιστής ασκεί έλεγχο, προσδιορίζοντας ποιο μέρος των δεδομένων του υποβάλλεται σε επεξεργασία σε κάθε νήμα.

Η απάντηση στο ερώτημα γιατί οι προγραμματιστές επέλεξαν αυτόν τον συγκεκριμένο οργανισμό δεν είναι ασήμαντη. Ένας λόγος είναι ότι ένα μπλοκ είναι εγγυημένο ότι θα εκτελεστείσε ένα πολυεπεξεργαστή συσκευής, αλλά ένας πολυεπεξεργαστής μπορεί να εκτελέσει πολλά διαφορετικά μπλοκ. Οι υπόλοιποι λόγοι θα γίνουν σαφείς αργότερα στο άρθρο.

Ένα μπλοκ εργασιών (threads) εκτελείται σε έναν πολυεπεξεργαστή σε μέρη ή ομάδες, που ονομάζονται warps. Το τρέχον μέγεθος στημόνι σε κάρτες γραφικών με υποστήριξη CUDA είναι 32 νήματα. Οι εργασίες μέσα στο warp pool εκτελούνται σε στυλ SIMD, π.χ. Όλα τα νήματα μέσα στο warp μπορούν να εκτελέσουν μόνο μία εντολή κάθε φορά.

Εδώ πρέπει να γίνει μια προειδοποίηση. Στις σύγχρονες αρχιτεκτονικές τη στιγμή της συγγραφής αυτού του άρθρου, ο αριθμός των επεξεργαστών μέσα σε έναν πολυεπεξεργαστή είναι 8, όχι 32. Από αυτό προκύπτει ότι δεν εκτελείται ολόκληρο το warp ταυτόχρονα, χωρίζεται σε 4 μέρη, τα οποία εκτελούνται διαδοχικά (καθώς οι επεξεργαστές είναι κλιμακωτές) .

Αλλά, πρώτον, οι προγραμματιστές CUDA δεν ρυθμίζουν αυστηρά το μέγεθος του στημονιού. Στα έργα τους αναφέρουν την παράμετρο μεγέθους warp και όχι τον αριθμό 32. Δεύτερον, από λογική άποψη, το warp είναι η ελάχιστη ένωση νημάτων για την οποία μπορούμε να πούμε ότι όλα τα νήματα μέσα σε αυτό εκτελούνται ταυτόχρονα - και στο Ταυτόχρονα δεν υπάρχουν υποθέσεις για τα υπόλοιπα δεν θα γίνει το σύστημα.

3.1.1. Διακλάδωση

Αμέσως προκύπτει το ερώτημα: εάν την ίδια χρονική στιγμή όλα τα νήματα μέσα σε ένα στημόνι εκτελούν την ίδια εντολή, τότε τι γίνεται με τους κλάδους; Άλλωστε, αν υπάρχει κλάδος στον κώδικα του προγράμματος, τότε οι οδηγίες θα είναι διαφορετικές. Εδώ χρησιμοποιείται μια τυπική λύση για προγραμματισμό SIMD (Εικόνα 6).

Ρύζι. 6. Οργάνωση υποκαταστήματος σε SIMD

Ας υποθέσουμε ότι έχουμε τον ακόλουθο κώδικα:

αν (συνέχεια)ΣΙ;

Στην περίπτωση του SISD (Single Instruction Single Data), εκτελούμε την εντολή A, ελέγχουμε τη συνθήκη και μετά εκτελούμε τις προτάσεις B και D (αν η συνθήκη είναι αληθής).

Ας έχουμε τώρα 10 νήματα που τρέχουν σε στυλ SIMD. Και στα 10 νήματα εκτελούμε την πρόταση Α, μετά ελέγχουμε τη συνθήκη συνθήκης και αποδεικνύεται ότι σε 9 από τα 10 νήματα είναι αληθές και σε ένα νήμα είναι ψευδές.

Είναι σαφές ότι δεν μπορούμε να εκκινήσουμε 9 νήματα για να εκτελέσουμε τον τελεστή Β και τα υπόλοιπα για να εκτελέσουμε τον τελεστή C, επειδή μόνο μία εντολή μπορεί να εκτελεστεί σε όλα τα νήματα ταυτόχρονα. Σε αυτήν την περίπτωση, πρέπει να κάνετε αυτό: πρώτα, "σκοτώνουμε" το αποσπασματικό νήμα έτσι ώστε να μην χαλάσει τα δεδομένα κανενός και εκτελούμε τα 9 υπόλοιπα νήματα. Στη συνέχεια «σκοτώνουμε» 9 νήματα που εκτέλεσαν τον τελεστή Β και περνάμε από ένα νήμα με τον τελεστή C. Μετά από αυτό, τα νήματα συγχωνεύονται ξανά και εκτελούν τον τελεστή D όλα ταυτόχρονα.

Το θλιβερό αποτέλεσμα είναι ότι όχι μόνο σπαταλούνται οι πόροι του επεξεργαστή για λείανση κενών ψηφίων στα διαχωρισμένα νήματα, αλλά το πολύ χειρότερο είναι ότι θα καταλήξουμε να εκτελέσουμε ΚΑΙ ΔΥΟ κλάδους.

Ωστόσο, δεν είναι όλα τόσο άσχημα όσο φαίνεται με την πρώτη ματιά. Ένα πολύ μεγάλο πλεονέκτημα της τεχνολογίας είναι ότι αυτά τα κόλπα εκτελούνται δυναμικά από τον οδηγό CUDA και είναι απολύτως διαφανή στον προγραμματιστή. Ταυτόχρονα, όταν ασχολούμαστε με εντολές SSE των σύγχρονων CPU (συγκεκριμένα στην περίπτωση της προσπάθειας εκτέλεσης 4 αντιγράφων του αλγορίθμου ταυτόχρονα), ο ίδιος ο προγραμματιστής πρέπει να φροντίσει τις λεπτομέρειες: συνδυάστε δεδομένα σε τετράγωνα, μην ξεχνάτε την ευθυγράμμιση , και γενικά γράφουν σε χαμηλό επίπεδο, μάλιστα, how in assembler.

Από όλα τα παραπάνω προκύπτει ένα πολύ σημαντικό συμπέρασμα. Τα κλαδιά δεν προκαλούν από μόνα τους υποβάθμιση της απόδοσης. Τα μόνα επιβλαβή κλαδιά είναι εκείνα όπου τα νήματα αποκλίνουν μέσα στην ίδια δεξαμενή νημάτων στημονιού. Επιπλέον, εάν τα νήματα αποκλίνουν μέσα σε ένα μπλοκ, αλλά σε διαφορετικές ομάδες στημονιού ή μέσα σε διαφορετικά μπλοκ, αυτό δεν έχει κανένα απολύτως αποτέλεσμα.

3.1.2. Επικοινωνία μεταξύ των νημάτων

Τη στιγμή της συγγραφής αυτού του άρθρου, οποιαδήποτε αλληλεπίδραση μεταξύ νημάτων (συγχρονισμός και ανταλλαγή δεδομένων) ήταν δυνατή μόνο εντός ενός μπλοκ. Δηλαδή, είναι αδύνατο να οργανωθεί η αλληλεπίδραση μεταξύ νημάτων διαφορετικών μπλοκ χρησιμοποιώντας μόνο τεκμηριωμένες δυνατότητες.

Όσον αφορά τα μη τεκμηριωμένα χαρακτηριστικά, η χρήση τους αποθαρρύνεται ιδιαίτερα. Ο λόγος για αυτό είναι ότι βασίζονται στα συγκεκριμένα χαρακτηριστικά υλικού ενός συγκεκριμένου συστήματος.

Ο συγχρονισμός όλων των εργασιών μέσα σε ένα μπλοκ πραγματοποιείται καλώντας τη συνάρτηση __synchtreads. Η ανταλλαγή δεδομένων είναι δυνατή μέσω της κοινής μνήμης, καθώς είναι κοινή για όλες τις εργασίες εντός του μπλοκ.

3.2. Μνήμη

Το CUDA διακρίνει έξι τύπους μνήμης (Εικ. 7). Πρόκειται για καταχωρητές, τοπική, καθολική, κοινόχρηστη, σταθερή και υφή μνήμη.

Αυτή η αφθονία οφείλεται στις ιδιαιτερότητες της κάρτας γραφικών και στον πρωταρχικό σκοπό της, καθώς και στην επιθυμία των προγραμματιστών να κάνουν το σύστημα όσο το δυνατόν φθηνότερο, θυσιάζοντας σε διάφορες περιπτώσεις είτε την ευελιξία είτε την ταχύτητα.

Ρύζι. 7. Τύποι μνήμης στο CUDA

3.2.0. Μητρώα

Όποτε είναι δυνατόν, ο μεταγλωττιστής προσπαθεί να τοποθετήσει όλες τις τοπικές μεταβλητές συναρτήσεων σε καταχωρητές. Τέτοιες μεταβλητές έχουν πρόσβαση με τη μέγιστη ταχύτητα. Στην τρέχουσα αρχιτεκτονική, υπάρχουν 8192 καταχωρητές 32-bit διαθέσιμοι ανά πολυεπεξεργαστή. Για να προσδιορίσετε πόσοι καταχωρητές είναι διαθέσιμοι σε ένα νήμα, πρέπει να διαιρέσετε αυτόν τον αριθμό (8192) με το μέγεθος του μπλοκ (τον αριθμό των νημάτων σε αυτό).

Με τη συνήθη διαίρεση των 64 νημάτων ανά μπλοκ, το αποτέλεσμα είναι μόνο 128 καταχωρητές (υπάρχουν ορισμένα αντικειμενικά κριτήρια, αλλά το 64 είναι κατάλληλο κατά μέσο όρο για πολλές εργασίες). Στην πραγματικότητα, το nvcc δεν θα εκχωρήσει ποτέ 128 καταχωρητές. Συνήθως δεν δίνει περισσότερες από 40 και οι υπόλοιπες μεταβλητές θα μπουν στην τοπική μνήμη. Αυτό συμβαίνει επειδή πολλά μπλοκ μπορούν να εκτελεστούν σε έναν πολυεπεξεργαστή. Ο μεταγλωττιστής προσπαθεί να μεγιστοποιήσει τον αριθμό των μπλοκ που λειτουργούν ταυτόχρονα. Για μεγαλύτερη αποτελεσματικότητα, θα πρέπει να προσπαθήσετε να καταλάβετε λιγότερους από 32 καταχωρητές. Στη συνέχεια, θεωρητικά, μπορούν να εκκινηθούν 4 μπλοκ (8 παραμορφώσεις, εάν 64 νήματα σε ένα μπλοκ) σε έναν πολυεπεξεργαστή. Ωστόσο, εδώ θα πρέπει επίσης να λάβετε υπόψη την ποσότητα της κοινόχρηστης μνήμης που καταλαμβάνεται από τα νήματα, καθώς εάν ένα μπλοκ καταλαμβάνει ολόκληρη την κοινόχρηστη μνήμη, δύο τέτοια μπλοκ δεν μπορούν να εκτελεστούν ταυτόχρονα σε έναν πολυεπεξεργαστή.

3.2.1. Τοπική μνήμη

Σε περιπτώσεις όπου τα δεδομένα τοπικής διαδικασίας είναι πολύ μεγάλα ή ο μεταγλωττιστής δεν μπορεί να υπολογίσει κάποιο σταθερό βήμα πρόσβασης για αυτά, μπορεί να τα τοποθετήσει στην τοπική μνήμη. Αυτό μπορεί να διευκολυνθεί, για παράδειγμα, με τη χύτευση δεικτών για τύπους διαφορετικών μεγεθών.

Φυσικά, η τοπική μνήμη είναι ανάλογη με την παγκόσμια μνήμη και λειτουργεί με την ίδια ταχύτητα. Τη στιγμή της σύνταξης, δεν υπάρχουν μηχανισμοί που να εμποδίζουν ρητά τον μεταγλωττιστή να χρησιμοποιεί τοπική μνήμη για συγκεκριμένες μεταβλητές. Εφόσον είναι αρκετά δύσκολος ο έλεγχος της τοπικής μνήμης, είναι προτιμότερο να μην τη χρησιμοποιήσετε καθόλου (βλ. ενότητα 4 «Προτάσεις βελτιστοποίησης»).

3.2.2. Καθολική μνήμη

Στην τεκμηρίωση CUDA ως ένα από τα κύρια επιτεύγματαΗ τεχνολογία παρέχει τη δυνατότητα αυθαίρετης διευθυνσιοδότησης της παγκόσμιας μνήμης. Δηλαδή, μπορείτε να διαβάσετε από οποιοδήποτε κελί μνήμης και μπορείτε επίσης να γράψετε σε ένα αυθαίρετο κελί (αυτό συνήθως δεν συμβαίνει σε μια GPU).

Ωστόσο, σε αυτή την περίπτωση θα πρέπει να πληρώσετε για την ευελιξία με ταχύτητα. Η παγκόσμια μνήμη δεν αποθηκεύεται προσωρινά. Λειτουργεί πολύ αργά, ο αριθμός των προσβάσεων στην παγκόσμια μνήμη θα πρέπει να ελαχιστοποιηθεί σε κάθε περίπτωση.

Η συνολική μνήμη απαιτείται κυρίως για την αποθήκευση των αποτελεσμάτων ενός προγράμματος πριν από την αποστολή τους στον κεντρικό υπολογιστή (σε κανονική μνήμη DRAM). Ο λόγος για αυτό είναι ότι η παγκόσμια μνήμη είναι το μόνο είδος μνήμης όπου μπορείτε να γράψετε οτιδήποτε.

Οι μεταβλητές που δηλώνονται με τον προσδιορισμό __global__ εκχωρούνται στην καθολική μνήμη. Η καθολική μνήμη μπορεί επίσης να εκχωρηθεί δυναμικά καλώντας cudaMalloc(void* mem, int size) στον κεντρικό υπολογιστή. Αυτή η λειτουργία δεν μπορεί να κληθεί από τη συσκευή. Ως εκ τούτου, η εκχώρηση μνήμης θα πρέπει να γίνεται από το κεντρικό πρόγραμμα που εκτελείται στη CPU. Τα δεδομένα από τον κεντρικό υπολογιστή μπορούν να σταλούν στη συσκευή καλώντας τη συνάρτηση cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Μπορείτε να κάνετε την αντίστροφη διαδικασία με τον ίδιο ακριβώς τρόπο:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Αυτή η κλήση γίνεται επίσης από τον οικοδεσπότη.

Όταν εργάζεστε με την παγκόσμια μνήμη, είναι σημαντικό να ακολουθείτε τον κανόνα της συγχώνευσης. Η κύρια ιδέα είναι ότι τα νήματα πρέπει να έχουν πρόσβαση σε διαδοχικά κελιά μνήμης, 4, 8 ή 16 byte. Σε αυτήν την περίπτωση, το πρώτο νήμα πρέπει να έχει πρόσβαση σε μια διεύθυνση ευθυγραμμισμένη με όριο 4, 8 ή 16 byte, αντίστοιχα. Οι διευθύνσεις που επιστρέφονται από το cudaMalloc είναι ευθυγραμμισμένες σε ένα όριο τουλάχιστον 256 byte.

3.2.3. Κοινή μνήμη

Η κοινόχρηστη μνήμη δεν είναι προσωρινή, αλλά γρήγορη. Συνιστάται η χρήση του ως διαχειριζόμενη κρυφή μνήμη. Μόνο 16 KB κοινόχρηστης μνήμης είναι διαθέσιμα ανά πολυεπεξεργαστή. Διαιρώντας αυτόν τον αριθμό με τον αριθμό των εργασιών στο μπλοκ, λαμβάνουμε τη μέγιστη διαθέσιμη κοινόχρηστη μνήμη ανά νήμα (αν σκοπεύετε να τη χρησιμοποιήσετε ανεξάρτητα σε όλα τα νήματα).

Ένα χαρακτηριστικό γνώρισμα της κοινόχρηστης μνήμης είναι ότι αντιμετωπίζεται εξίσου για όλες τις εργασίες μέσα σε ένα μπλοκ (Εικ. 7). Συνεπάγεται ότι μπορεί να χρησιμοποιηθεί για την ανταλλαγή δεδομένων μεταξύ νημάτων ενός μόνο μπλοκ.

Είναι εγγυημένο ότι κατά την εκτέλεση μπλοκ σε έναν πολυεπεξεργαστή, τα περιεχόμενα της κοινόχρηστης μνήμης θα διατηρηθούν. Ωστόσο, μετά την αλλαγή ενός μπλοκ σε έναν πολυεπεξεργαστή, δεν είναι εγγυημένο ότι τα περιεχόμενα του παλιού μπλοκ θα διατηρηθούν. Επομένως, δεν πρέπει να προσπαθήσετε να συγχρονίσετε εργασίες μεταξύ μπλοκ, αφήνοντας δεδομένα στην κοινόχρηστη μνήμη και ελπίζοντας για την ασφάλειά τους.

Οι μεταβλητές που δηλώνονται με τον προσδιορισμό __shared__ εκχωρούνται στην κοινόχρηστη μνήμη.

Shared__ float mem_shared;

Θα πρέπει να τονιστεί για άλλη μια φορά ότι υπάρχει μόνο μία κοινόχρηστη μνήμη για ένα μπλοκ. Επομένως, εάν πρέπει να το χρησιμοποιήσετε απλώς ως διαχειριζόμενη κρυφή μνήμη, θα πρέπει να έχετε πρόσβαση σε διαφορετικά στοιχεία του πίνακα, για παράδειγμα, όπως αυτό:

float x = mem_shared;

Όπου threadIdx.x είναι ο δείκτης x του νήματος μέσα στο μπλοκ.

3.2.4. Σταθερή μνήμη

Η σταθερή μνήμη αποθηκεύεται στην κρυφή μνήμη, όπως φαίνεται στην Εικ. 4. Η κρυφή μνήμη υπάρχει σε ένα μόνο αντίγραφο για έναν πολυεπεξεργαστή, πράγμα που σημαίνει ότι είναι κοινή για όλες τις εργασίες εντός του μπλοκ. Στον κεντρικό υπολογιστή, μπορείτε να γράψετε κάτι στη σταθερή μνήμη καλώντας τη συνάρτηση cudaMemcpyToSymbol. Από τη συσκευή, η σταθερή μνήμη είναι μόνο για ανάγνωση.

Η σταθερή μνήμη είναι πολύ βολική στη χρήση. Μπορείτε να τοποθετήσετε δεδομένα οποιουδήποτε τύπου σε αυτό και να τα διαβάσετε χρησιμοποιώντας μια απλή ανάθεση.

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ σημαίνει ότι το device_kernel είναι ένας πυρήνας που μπορεί να εκτελεστεί στη GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ΛΑΘΟΣ! Η σταθερή μνήμη είναι μόνο για ανάγνωση

Δεδομένου ότι η σταθερή μνήμη χρησιμοποιεί μια προσωρινή μνήμη, η πρόσβαση σε αυτήν είναι γενικά αρκετά γρήγορη. Το μόνο αλλά πολύ μεγάλο μειονέκτημα της σταθερής μνήμης είναι ότι το μέγεθός της είναι μόνο 64 KB (για ολόκληρη τη συσκευή). Αυτό υποδηλώνει ότι είναι λογικό να αποθηκεύεται μόνο μια μικρή ποσότητα δεδομένων που χρησιμοποιούνται συχνά στη μνήμη περιβάλλοντος.

3.2.5. Μνήμη υφής

Η μνήμη υφής αποθηκεύεται στην κρυφή μνήμη (Εικ. 4). Υπάρχει μόνο μία κρυφή μνήμη για κάθε πολυεπεξεργαστή, πράγμα που σημαίνει ότι αυτή η κρυφή μνήμη είναι κοινή για όλες τις εργασίες εντός του μπλοκ.

Το όνομα της μνήμης υφής (και, δυστυχώς, της λειτουργικότητας) κληρονομείται από τις έννοιες "υφή" και "υφή". Η υφή είναι η διαδικασία εφαρμογής μιας υφής (απλώς μιας εικόνας) σε ένα πολύγωνο κατά τη διαδικασία ραστεροποίησης. Η μνήμη υφής είναι βελτιστοποιημένη για 2D δειγματοληψία δεδομένων και έχει τις ακόλουθες δυνατότητες:

γρήγορη ανάκτηση τιμών σταθερού μεγέθους (byte, λέξη, διπλή ή τετραπλή λέξη) από έναν μονοδιάστατο ή δισδιάστατο πίνακα.

κανονικοποιημένη διευθυνσιοδότηση με αριθμούς float στο διάστημα . Στη συνέχεια, μπορείτε να τα επιλέξετε χρησιμοποιώντας κανονικοποιημένη διευθυνσιοδότηση. Η τιμή που προκύπτει θα είναι μια λέξη του τύπου float4 που αντιστοιχίζεται στο διάστημα ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); //κατανομή μνήμης στη GPU

// ρύθμιση παραμέτρων υφής υφή

Texture.addressMode = cudaAddressModeWrap; //τρόπος Κάλυμμα

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //πλησιέστερη τιμή

Texture.normalized = false; // μην χρησιμοποιείτε κανονικοποιημένη διευθυνσιοδότηση

CudaBindTexture(0, texture, gpu_memory, N ) // από εδώ και στο εξής αυτή η μνήμη θα θεωρείται μνήμη υφής


CudaMemcpy(gpu_memory, cpu_buffer, N*sizeof(uint 4), cudaMemcpyHostToDevice ); // αντιγραφή δεδομένων σεGPU

// __global__ σημαίνει ότι το device_kernel είναι ο πυρήνας που πρέπει να παραλληλιστεί

Global__ void device_kernel()

uint4 a = tex1Dfetch(υφή,0); // μπορείτε να επιλέξετε δεδομένα μόνο με αυτόν τον τρόπο!

uint4 b = tex1Dfetch(υφή,1);

int c = a.x * b.y;

...

3.3. Απλό παράδειγμα

Ως απλό παράδειγμα, εξετάστε το πρόγραμμα cppIntegration από το CUDA SDK. Επιδεικνύει τεχνικές εργασίας με CUDA, καθώς και τη χρήση του nvcc (ένα ειδικό υποσύνολο του μεταγλωττιστή C++ από τη Nvidia) σε συνδυασμό με το MS Visual Studio, το οποίο απλοποιεί σημαντικά την ανάπτυξη προγραμμάτων στο CUDA.

4.1. Αναλύστε σωστά την εργασία σας

Δεν είναι όλες οι εργασίες κατάλληλες για αρχιτεκτονικές SIMD. Εάν η εργασία σας δεν είναι κατάλληλη για αυτό, μπορεί να μην αξίζει να χρησιμοποιήσετε μια GPU. Αλλά αν είστε αποφασισμένοι να χρησιμοποιήσετε μια GPU, θα πρέπει να προσπαθήσετε να χωρίσετε τον αλγόριθμο σε κομμάτια που μπορούν να εκτελεστούν αποτελεσματικά σε στυλ SIMD. Εάν είναι απαραίτητο, αλλάξτε τον αλγόριθμο για να λύσετε το πρόβλημά σας, βρείτε έναν νέο - έναν που θα ταίριαζε καλά με το SIMD. Ένα παράδειγμα κατάλληλης περιοχής για τη χρήση της GPU είναι η υλοποίηση της πυραμιδικής προσθήκης στοιχείων πίνακα.

4.2. Επιλογή του τύπου μνήμης

Τοποθετήστε τα δεδομένα σας σε υφή ή σταθερή μνήμη εάν όλες οι εργασίες στο ίδιο μπλοκ έχουν πρόσβαση στην ίδια θέση μνήμης ή σε κοντινές περιοχές. Τα δεδομένα 2D μπορούν να υποστούν αποτελεσματική επεξεργασία χρησιμοποιώντας τις λειτουργίες text2Dfetch και text2D. Η μνήμη υφής είναι ειδικά βελτιστοποιημένη για δειγματοληψία 2D.

Χρησιμοποιήστε την καθολική μνήμη σε συνδυασμό με την κοινόχρηστη μνήμη εάν όλες οι εργασίες έχουν τυχαία πρόσβαση σε διαφορετικές, ευρέως διαχωρισμένες τοποθεσίες μνήμης (με πολύ διαφορετικές διευθύνσεις ή συντεταγμένες εάν πρόκειται για δεδομένα 2D/3D).

παγκόσμια μνήμη => κοινόχρηστη μνήμη

Syncthreads();

Επεξεργασία δεδομένων σε κοινόχρηστη μνήμη

Syncthreads();

καθολική μνήμη<= разделяемая память

4.3. Ενεργοποίηση μετρητών μνήμης

Η σημαία μεταγλωττιστή --ptxas-options=-v σάς επιτρέπει να πείτε ακριβώς πόση και τι είδους μνήμη (μητρώα, κοινόχρηστη, τοπική, σταθερή) χρησιμοποιείτε. Εάν ο μεταγλωττιστής χρησιμοποιεί τοπική μνήμη, σίγουρα το γνωρίζετε. Η ανάλυση δεδομένων σχετικά με την ποσότητα και τον τύπο της μνήμης που χρησιμοποιείται μπορεί να σας βοηθήσει πολύ στη βελτιστοποίηση του προγράμματός σας.

4.4. Προσπαθήστε να ελαχιστοποιήσετε τη χρήση καταχωρητών και κοινόχρηστης μνήμης

Όσο περισσότερους καταχωρητές ή κοινόχρηστη μνήμη χρησιμοποιεί ο πυρήνας, τόσο λιγότερα νήματα (ή μάλλον παραμορφώσεις) μπορούν να εκτελεστούν ταυτόχρονα σε έναν πολυεπεξεργαστή, επειδή Οι πόροι πολλαπλών επεξεργαστών είναι περιορισμένοι. Επομένως, μια ελαφρά αύξηση της πληρότητας των καταχωρητών ή της κοινής μνήμης μπορεί σε ορισμένες περιπτώσεις να οδηγήσει σε πτώση της απόδοσης στο μισό - ακριβώς επειδή τώρα εκτελούνται ακριβώς τα μισά warps ταυτόχρονα σε έναν πολυεπεξεργαστή.

4.5. Κοινόχρηστη μνήμη αντί για τοπική

Εάν ο μεταγλωττιστής Nvidia για κάποιο λόγο εκχώρησε δεδομένα στην τοπική μνήμη (συνήθως αυτό είναι αισθητό από μια πολύ μεγάλη πτώση της απόδοσης σε μέρη όπου δεν υπάρχει τίποτα με ένταση πόρων), μάθετε ακριβώς ποια δεδομένα κατέληξαν στην τοπική μνήμη και τοποθετήστε τα σε κοινόχρηστο μνήμη).

Συχνά ο μεταγλωττιστής τοποθετεί μια μεταβλητή στην τοπική μνήμη εάν δεν χρησιμοποιείται συχνά. Για παράδειγμα, αυτό είναι κάποιο είδος συσσωρευτή όπου συσσωρεύετε μια τιμή υπολογίζοντας κάτι σε έναν βρόχο. Εάν ο βρόχος είναι μεγάλος ως προς τον όγκο του κώδικα (αλλά όχι στον χρόνο εκτέλεσης!), τότε ο μεταγλωττιστής μπορεί να τοποθετήσει τον συσσωρευτή σας στην τοπική μνήμη, επειδή χρησιμοποιείται σχετικά σπάνια και τα μητρώα είναι λίγα. Η απώλεια απόδοσης σε αυτή την περίπτωση μπορεί να είναι αισθητή.

Εάν πραγματικά χρησιμοποιείτε σπάνια μια μεταβλητή, είναι προτιμότερο να την τοποθετήσετε ρητά στην καθολική μνήμη.

Αν και μπορεί να φαίνεται βολικό για τον μεταγλωττιστή να τοποθετεί αυτόματα τέτοιες μεταβλητές στην τοπική μνήμη, στην πραγματικότητα δεν είναι έτσι. Θα είναι δύσκολο να βρεθεί το σημείο συμφόρησης κατά τις επόμενες τροποποιήσεις του προγράμματος, εάν η μεταβλητή αρχίσει να χρησιμοποιείται πιο συχνά. Ο μεταγλωττιστής μπορεί ή όχι να μεταφέρει μια τέτοια μεταβλητή για να καταχωρήσει τη μνήμη. Εάν ο τροποποιητής __global__ καθορίζεται ρητά, ο προγραμματιστής είναι πιο πιθανό να του δώσει προσοχή.

4.6. Ξετυλίγοντας βρόχους

Το ξετύλιγμα βρόχου είναι μια τυπική τεχνική απόδοσης σε πολλά συστήματα. Η ουσία του είναι να εκτελεί περισσότερες ενέργειες σε κάθε επανάληψη, μειώνοντας έτσι τον συνολικό αριθμό των επαναλήψεων, και συνεπώς τον αριθμό των διακλαδώσεων υπό όρους που θα πρέπει να εκτελέσει ο επεξεργαστής.

Δείτε πώς μπορείτε να ξετυλίξετε τον βρόχο για την εύρεση του αθροίσματος ενός πίνακα (για παράδειγμα, ενός ακέραιου αριθμού):

int a[N]; int άθροισμα?

για (int i=0;i

Φυσικά, οι βρόχοι μπορούν επίσης να ξετυλιχθούν χειροκίνητα (όπως φαίνεται παραπάνω), αλλά αυτό είναι μη παραγωγική δουλειά. Είναι πολύ καλύτερο να χρησιμοποιείτε πρότυπα C++ σε συνδυασμό με ενσωματωμένες συναρτήσεις.

πρότυπο

κλάση ArraySumm

Device__ static T exec(const T* arr) ( return arr + ArraySumm (arr+1); )

πρότυπο

κλάση ArraySumm<0,T>

Device__ static T exec(const T* arr) (επιστροφή 0; )

για (int i=0;i

summ+= ArraySumm<4,int>::exec(a);

Πρέπει να σημειωθεί ένα ενδιαφέρον χαρακτηριστικό του μεταγλωττιστή nvcc. Ο μεταγλωττιστής θα ενσωματώνει πάντα συναρτήσεις τύπου __συσκευή__ από προεπιλογή (υπάρχει μια ειδική οδηγία __noinline__ για να την παρακάμψει).

Επομένως, μπορείτε να είστε βέβαιοι ότι ένα παράδειγμα όπως το παραπάνω θα ξεδιπλωθεί σε μια απλή ακολουθία δηλώσεων και δεν θα είναι σε καμία περίπτωση κατώτερο σε αποτελεσματικότητα από τον χειρόγραφο κώδικα. Ωστόσο, στη γενική περίπτωση (όχι nvcc) δεν μπορείτε να είστε σίγουροι για αυτό, καθώς το inline είναι μόνο μια ένδειξη στον μεταγλωττιστή ότι μπορεί να αγνοήσει. Επομένως, δεν είναι εγγυημένο ότι οι λειτουργίες σας θα είναι ενσωματωμένες.

4.7. Στοίχιση δεδομένων και δειγματοληψία 16 byte

Ευθυγραμμίστε τις δομές δεδομένων σε όρια 16 byte. Σε αυτήν την περίπτωση, ο μεταγλωττιστής θα μπορεί να χρησιμοποιήσει ειδικές οδηγίες για αυτούς που φορτώνουν δεδομένα 16 byte τη φορά.

Εάν η δομή είναι 8 byte ή λιγότερο, μπορείτε να την ευθυγραμμίσετε στα 8 byte. Αλλά σε αυτήν την περίπτωση, μπορείτε να επιλέξετε δύο μεταβλητές ταυτόχρονα συνδυάζοντας δύο μεταβλητές 8 byte σε μια δομή χρησιμοποιώντας μια ένωση ή cast δείκτη. Το Casting θα πρέπει να χρησιμοποιείται με προσοχή γιατί ο μεταγλωττιστής μπορεί να τοποθετήσει τα δεδομένα στην τοπική μνήμη και όχι σε καταχωρητές.

4.8. Συγκρούσεις τράπεζας κοινής μνήμης

Η κοινόχρηστη μνήμη είναι οργανωμένη με τη μορφή 16 (συνολικά!) τραπεζών μνήμης με βήμα 4 byte. Κατά την εκτέλεση μιας ομάδας νημάτων στημόνι σε έναν πολυεπεξεργαστή, χωρίζεται σε δύο μισά (αν warp-size = 32) των 16 νημάτων, τα οποία έχουν πρόσβαση στην κοινόχρηστη μνήμη με τη σειρά τους.

Οι εργασίες στα διαφορετικά μισά του στημονιού δεν έρχονται σε αντίθεση με την κοινόχρηστη μνήμη. Λόγω του γεγονότος ότι οι εργασίες του μισού του warp pool θα έχουν πρόσβαση στις ίδιες τράπεζες μνήμης, θα προκύψουν συγκρούσεις και, ως αποτέλεσμα, μείωση της απόδοσης. Οι εργασίες εντός του μισού του warp μπορούν να έχουν πρόσβαση σε διαφορετικές περιοχές της κοινόχρηστης μνήμης σε ένα συγκεκριμένο βήμα.

Τα βέλτιστα βήματα είναι 4, 12, 28, ..., 2^n-4 byte (Εικ. 8).



Ρύζι. 8. Βέλτιστα βήματα.

Τα μη βέλτιστα βήματα είναι 1, 8, 16, 32, ..., 2^n byte (Εικ. 9).



Ρύζι. 9. Υποβέλτιστα βήματα

4.9. Ελαχιστοποίηση των κινήσεων δεδομένων κεντρικού υπολογιστή<=>Συσκευή

Προσπαθήστε να μεταφέρετε τα ενδιάμεσα αποτελέσματα στον κεντρικό υπολογιστή για επεξεργασία χρησιμοποιώντας την CPU όσο το δυνατόν λιγότερο. Εφαρμόστε, αν όχι ολόκληρο τον αλγόριθμο, τουλάχιστον το κύριο μέρος του στη GPU, αφήνοντας μόνο τις εργασίες ελέγχου στην CPU.

5. Φορητή βιβλιοθήκη μαθηματικών CPU/GPU

Ο συγγραφέας αυτού του άρθρου έχει γράψει μια φορητή βιβλιοθήκη MGML_MATH για εργασία με απλά χωροαντικείμενα, ο κώδικας της οποίας είναι λειτουργικός τόσο στη συσκευή όσο και στον κεντρικό υπολογιστή.

Η βιβλιοθήκη MGML_MATH μπορεί να χρησιμοποιηθεί ως πλαίσιο για τη συγγραφή φορητών (ή υβριδικών) συστημάτων CPU/GPU για τον υπολογισμό φυσικών, γραφικών ή άλλων χωρικών προβλημάτων. Το κύριο πλεονέκτημά του είναι ότι ο ίδιος κώδικας μπορεί να χρησιμοποιηθεί τόσο στη CPU όσο και στη GPU, και ταυτόχρονα, η ταχύτητα είναι η κύρια απαίτηση για τη βιβλιοθήκη.

6 . Βιβλιογραφία

Κρις Κάσπερσκι. Τεχνικές βελτιστοποίησης προγραμμάτων. Αποτελεσματική χρήση της μνήμης. - Αγία Πετρούπολη: BHV-Petersburg, 2003. - 464 σελ.: ill.

Οδηγός προγραμματισμού CUDA 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

Οδηγός προγραμματισμού CUDA 1.1. σελίδα 14-15

Οδηγός προγραμματισμού CUDA 1.1. σελίδα 48