THMMY.gr

Topic που αφορά την 3η εργασία του μαθήματοΚαλοφάγωτη.

Μήπως κατάλαβε κανείς τι αντιπροσώπευει η παράμετρος p = [0.33 0.45 0.66] που αναφέρει η εκφώνηση ;;

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

Α μαλιστα φίλε.  Και την ρουτίνα ματλαμπ που μας δίνεται θα την μετατρέψουμε σε c κώδικα εμείς , ή γίνεται να την δώσουμε αυτόματα ως όρισμα μέσω ματλαμπ στο c πρόγραμμα μας ;;

+1

Όντως γιατί μας το έδωσε σε Matlab ?

η ρουτινα στο ματλαμπ απ'οτι καταλαβαινω ειναι για να παραγουμε τυχαια εισοδο για το προγραμμα μας. τωρα ειτε θα το παιρνουμε κατ'ευθειαν απο το ματλαμπ, ειτε ξερω γω κοπυ πειστ? δεν ξερω. αλλα και η εξαρχης υλοποιηση σε c αυτης της ρουτινας δεν ειναι κακη ιδεα. το p ειναι παραμετρος που εχει στην ρουτινα, ποσο συχνα να μην βαζει τιμη "απειρο" στον γραφο, δηλαδη ποσο συχνα ενωνονται δυο σημεια

δειτε και αυτo:
http://www.cs.ucsb.edu/~gilbert/cs140/old/cs140Win2011/cudaproject/apsp.html

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

Γιατί να το γράψετε σε C; Υποθέτω ότι μας την έδωσε για να έχουν όλοι τις ίδιες εισόδους και να μην ασχοληθούμε τόσο με αυτό

Για να μην εγκαθιστάς matlab σε linux μόνο και μόνο για να τρέχεις μία συνάρτηση  :D

-Octave

Εκτός απο αυτό, πως ακριβώς θα μεταφέρουμε τον πίνακα απο Matlab σε C ?

Αρχεία;

αρχεια? μπελας! καλυτερα κατευθειαν σε c και να παει στο διαολο

Ρε παιδιά όταν λέει "Ένα κελί ανα νήμα" εννοεί ότι κάθε block θα έχει 1 νήμα(thread);
Και όταν λέει "Πολλαπλά κελιά ανα νήμα" εννοεί ότι κάθε block θα έχει παραπάνω απο ένα νήμα(thread)?

Παιδιά, πως κανουμε compile cuda στον διαδη; Προσπαθω να κανω compile ενα πολυ απλο προγραμμα που προσθετει διανυσματα, και μου βγαζει ΑΠΕΙΡΑ ερρορ, οτι λειπουν παρενθεσεις, ; , οτι μεταβλητες δεν ειναι δηλωμενες και ενα καρο αλλα πραγματα, ενω το αρχειο ειναι 30 γραμμες το πηρα ετοιμο και δε φαινεται να εχει θεμα... Αποτι εχω καταλαβει πρεπει να γραψω:

nvcc myCuda.cu

Μηπως δεν ειναι αυτο; Πρεπει να κανουμε include καποια βιβλιοθηκη απο καπου (οπως στο Cilk);

+1. Αυτο το ενα κελι ανα νημα, μου φαινεται περιεργο. Εννοει ενα κελι του πινακα που θα εχουμε η κατι αλλο;

Δεν εχω διαβασει την εκφωνηση,αλλα γενικα αμα εχετε να κανετε με πινακα, τοτε καθε νημα αναλαμβανει ενα κελι του πινακα. Και στα πολλαπλα ανα νημα,πολλαπλα κελια  :D

Απλά το ανεβάζω.. μου φάνηκε ενδιαφέρον...
https://users.ics.aalto.fi/suomela/apsp/ (https://users.ics.aalto.fi/suomela/apsp/)

Yποθέτω πως θέλει να πει (threads per block) νήματα ανά κελί ;;;;

To κελιά ανά νήμα δεν μπορώ να το καταλάβω...

παίδες, έχουμε καμιά ιδέα για το πόσα μπλοκ μπορεί να σηκώσει ταυτόχρονα ο διαδης ?
γιατι στο πρώτο κέρνελ, όλα τρέχουν πόπα και είναι όλα σωστά με καλούς χρόνους εκτός απο 2^12....

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

Προς το παρόν έχω δοκιμάσει για 256 και για 512 θρεντς ανα μπλοκ. για 512 θρεντς ανα μπλοκ χρειάζονται συνολικά 32768 μπλοκ έτσι ώστε να έχουμε ενα κελί ανα νήμα.
Καμιά ιδέα ?

Το βασικό είναι να μην ξεπεράσεις τα Threads per block που δικαιούσαι από την εκάστοτε κάρτα γραφικών.Επειδή βλέπω πως έχεις πρόβλημα στα "πολλά"¨στοιχεία ,
προσπάθησε να κρατήσεις σταθερό τον αριθμό των threads λογικά στα 512 ή στα 1024 και δοκίμασε να δεις.Τα μπλοκς δεν μου φαίνονται πολλά για να έχεις θέμα...

Άκυρο τελικά, είχα ξεχάσει κάπου ενα 256 ενω έπρεπε να ήταν BLOCK_SIZE, οπότε στα πολλά στοιχεία έφτιαχνε παραπανω μπλοκ... :)

ανάφερε τίποτα για παράταση; γιατί με εξεταστική μπροστά μας, πολύ ζόρικα μέχρι τις 24 να βγει η εργασια

ειπε οτι θα παει μεχρι μια εβδομαδα μετα την εξεταστικη, κ οτι μπορουμε μεχρι τοτε να στειλουμε κ οτι διορθωσεις εχουμε για την δευτερη εργασια

Άκυρο! Μπέρδεψα τα κελιά του πίνακα με τα μπλοκς ! OK!

Για όσους τρέχουν την εργασία σε προσωπικό υπολογιστή, και όχι στον διαδη,
Παρακαλώ να προσέξουν το παρακάτω link ..  8))
 http://www.techworm.net/2016/01/nvidia-points-fingers-apple-google-chrome-incognito-mode-leaking-bug.html   (http://www.techworm.net/2016/01/nvidia-points-fingers-apple-google-chrome-incognito-mode-leaking-bug.html)
 :D :D :D

Αν έχει νιώσει κάποιος καλά την λογική της παραλληλοποίησης σε cuda, ας βοηθήσει λίγο με κάποιες απορίες.

- Με ποιον τρόπο θα χειριστούμε το k, της πρώτης λούπας για να βγάλουμε τα indexes? Μήπως είναι καλύτερα να καλούμε μέσω μιας επανάληψης εξωτερικά,  την συνάρτηση  που θα τρέξει στο device με δείκτη k ???  Κάπου πρέπει να πήρε το μάτι μου κάποια υλοποίηση όπου το k κάπως το ενσωμάτωνε εσωτερικά στην συνάρτηση του device. Μου φαίνεται σε κάποιο papper, αλλά μου φάνηκε κάπως σύνθετη. Ποιος είναι ο πιο απλό τρόπος.

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

-Με ποια λογική ορίζουμε το μέγεθος των μπλοκς?? ειδικά στο κομμάτι β, που ζητάει να χρησιμοποιηθεί η shared memory. Πως πρέπει να μοιράσουμε τα κελιά του πίνακα γράφου στα διάφορα blocks, ώστε κάθε μπλοκ να είναι ανεξάρτητο (όσο το δυνατό γίνεται) από τα υπόλοιπα και να μπορεί να χρησιμοποιεί για του τους υπολογισμούς την shared memory αντί της global? (Λογικά δεν επαρκεί μόνο η shared μνήμη, θα πρέπει να κάνουμε και κάποιες προσπελάσεις στην global κατά την διάρκεια των υπολογισμών, έτσι????)

-Τέλος όσον αφορά την αντιστοίχιση των κελιών στα νήματα των blocks. Έστω, ότι έχουμε φορτώσει τα στοιχεία του n*n πίνακα σε έναν πίνακα γραμμή για να τον φορτώσουμε στην συσκευή. Τότε, το πρώτο στοιχείο του πίνακα γραμμή αντιστοιχίζεται στο πρώτο νήμα του πρώτου μπλοκ, το στοιχείο 512 (μέγεθος block 256) αντιστοιχίζεται στο νήμα του 3ου μπλοκ έτσι ?? άρα μπορούμε να χρησιμοποιήσουμε για αρίθμηση το index=blockIdx.x * blockDim.x + threadIdx.x, ώστε να αναφερθούμε στην θέση του αρχικού πίνακα??? Και τη διαφορά έχει το παραπάνω index από αυτό Index = blockIdx.y * blockDim.y + threadIdx.y???

 Σορρυ για το κατεβατό…

Παιδιά πως κάνετε compile CUDA στο Διάδη?
μου πετάει αυτό το error ότι αρχείο και να κάνω compile...

/usr/bin/ld: final link failed: No space left on device
collect2: ld returned 1 exit status

δεν φταις εσυ, το βγαζει αρκετα συχνα

κ συνηθως στρωνει μετα απο καποιες ωρες

Σβήστε αρχεια απο τον server, εχει γεμίσει (παλι) ο δίσκος

ρε κηπουριδη τι στο καλο τρεχεις εδω και μια μερα?

Έστειλα μειλ και στον Πιτσιάνη και μου είπε ότι θα το ρυθμίσουν. Και να σβήσουμε και μεις ότι μπορούμε.

Παιδιά έχω αυτό το error:

In function `_start': (.text+0x20): undefined reference to `main'collect2: ld returned 1 exit status cuda error

Ξέρει κανένας τι παίζει. Για κάποιο λόγο δεν βρίσκει την main. Την έχω στο ίδιο αρχείο *.cu όπως και όλο το πρόγραμμα.

Kαλησπέρα , στο ζητούμενο 3 πρέπει να βρούμε καλύτερους ή χειρότερους χρόνους από το 2;

Προσωπικα σε εμενα ηταν αρκετα καλυτερο το 3, αλλα ισως αν εχεις κανει ομορφα το 2 να μην ισχυει απαραιτητα αυτο.
Προσπαθα παντως για καλυτερα.

Από e-thmmy:

Φλώρος

Δίνεται παράταση στην 3η εργασία.
Νέα ημερομηνία παράδοσης: 29/02

Εγώ νόμιζα ότι ούτως ή άλλως ίσχυε xD

Τυπικότητες....  :D

Τι χρονο σας κανει περιπου?

Για n = 12, p=1:

Ένα κελί ανά νήμα, χωρίς shared: 5.351042

Ένα κελί ανά νήμα, με shared: 4.904873

Σειριακός:  564.096380

(Στον Διάδη όλα αυτά.)

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

Edit:
Με βάση την συμβουλή από κάτω (thanks) έκανα μια ενδεικτική υλοποίηση.

Πολλά κελιά ανά νήμα, με shared: 3.503272

Δοκιμασε το πιο απλο πραγμα που σου ερχεται, μπορει να ειναι βελτιωμενο. Επειδη ειναι λιγο υπερβολη να καλεις μια καινουρια διεργασια για να αναλαβει μονο ενα κελι, μπορει να βελτιωθει και ετσι ξερα, παρολο που ακουγεται ηλιθιο.

Μια ακομα παρατηρηση, αν οι πυρηνες σου αναλαβουν πχ μια γραμμη με 8 κελια και ειναι δυο, ΜΗΝ κανει ο ενας τα 1-4 κι ο αλλος τα 5-8, αλλα καλυτερα ο ενας τα μονα κι ο αλλος τα ζυγα.
Επειδη οταν ζητας μια θεση μνημης, αυτοματα παιρνει και τις γειτονικες το συστημα με την υποθεση οτι θα χρειαστουν. Οποτε στη δευτερη μεθοδο το εκμεταλλευεσαι αυτο, στην πρωτη οχι.

Κατι ακομη, για τον ελεγχο ορθοτητας τα συγκρινετε με το αποτελεσμα που βγαζει ο σειριακος ή με τον αρχικο πινακα που δεν εχει υποστει καμια αλλαγη? Γιατι σε εμενα οι παραλληλες υλοποιησεις μου βγαζουν διαφορετικα αποτελεσματα απο τον σειριακο αλλα ειναι καλυτερες απο τον αρχικο πινακα Α

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

Για την υλοποιηση χωρις shared memory χρησιμοποιω αυτο το κομματι κωδικα...
 
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
 
if(row < N && col < N){
  for(k = 0; k < N; k++){
    if( A[N*row + col] > A[N*row + k] + A[k*N + col] )
      A[N*row + col] = A[N*row + k] + A[k*N + col];
    __syncthreads();
  }
}

Βλεπετε κανα λαθος? Δε μπορω να βγαλω ακρη γιατι δε λειτουργει οπως πρεπει...

Μιας και αρχιζω ΤΩΡΑ την εργασια :P :P :P ,
για πειτε λιγο. Τελικα δεν υπαρχει θεμα να κανουμε σε υλοποιηση c τον adjacency πινακα; ή πρέπει να μπλέξουμε με αρχεία;

Επίσης, ποιος ειναι ο ελεγχος ορθοτητας για τον σειριακο αλγοριθμο;

εγω παντως με c το εκανα... δε πιστευω να υπαρχει θεμα
Για τον ελεγχο ορθοτητας νομιζω οτι αν συγκρινεις τον πινακα που βρηκες με τον αρχικο και βρεις οτι τα κελια του ειναι μικροτερα ή ισα, αλλα ΟΧΙ ΟΛΑ ισα, τοτε εισαι οκ... Απ'την αλλη δε παιρνω και ορκο οτι αυτος ειναι και ο πιο ορθος τροπος. Αν καποιος εχει σκεφτει κατι καλυτερο is welcome.

Αν και υπάρχει περίπτωση ο γράφος να είναι ήδη ελάχιστος, πρακτικά αυτό που λέει ο nikos1 καλό είναι.

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

εγινε σας ευχαριστω!

Εχει κανεις ιδεα πως μπορουμε να τρεξουμε matlab scriptακι στον diades?

Ρε παίδες, την λούπα για το k την έχετε μέσα στο Kernel ή απ έξω ?

Αρχικά την είχα απ έξω (δλδ στην main και απλα καλούσα το kernel n-φορές για κάθε k)
και τώρα που προσπαθώ να το βάλω μέσα όλο μου βγάζει test failed...

Τίποτα ?

Απέξω, διότι πρέπει με κάποιο τρόπο να συγχρονίζεις όλα τα threads του πίνακα και όχι μόνο των blocks.

+1

Απλώς όταν είναι απ' έξω κάπως "δεν κανει ιδιαίτερη δουλειά" η χρήση της Shared..

π.χ.

με αυτό το κομμάτι κώδικα :


// Shared memory
   __shared__ float arr[3 * BLOCK_SIZE];
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   int i,j;

if(idx < n*n){
      i = idx / n;
      j = idx - i*n;

      arr[threadIdx.x] = dev_dist[i*n + j];
      arr[blockDim.x + threadIdx.x] = dev_dist[i*n + k];
      arr[2*blockDim.x + threadIdx.x] = dev_dist[k*n + j];
      

      if(arr[threadIdx.x] > (arr[blockDim.x + threadIdx.x] + arr[2*blockDim.x + threadIdx.x])){
         arr[threadIdx.x] = arr[blockDim.x + threadIdx.x] + arr[2*blockDim.x + threadIdx.x];
      }
      dev_dist[idx] = arr[threadIdx.x];
   }
   __syncthreads();


που χρησιμοποιώ, νιώθω οτι δε κάνει αρκετή δουλειά. Μέχρι n = 10 υπάρχει όντως βελτίωση στο χρόνο πρώτου κερνελ - δεύτερου κερνελ
Αλλά απο κει και πέρα το δεύτερο Kernel παράγει χειρότερους χρόνους απο το πρώτο! (καλα, οχι και πολυ χειρότερους, αλλά και πάλι .. )

Στις τελικές μετρήσεις που κάνατε είδατε καμιά ουσιαστική διαφορά για τις διάφορες τιμές της πιθανότητας p=[0.33  0.45  0.66]???

Σε εμένα τα αποτελέσματα για μια συγκεκριμένη τιμή του N βγαίνουν πάνω κάτω τα ίδια, για τις διάφορες τιμές του p.

Λογικό να υπάρχει κάποια πολύ μικρή διαφορά αφού για p=0.33 ο πίνακας περιέχει πιο πολλές INF τιμές σε σχέση με p=0.66, οπότε στην πρώτη περίπτωση πιθανόν θα γίνουν και περισσότερες ανανεώσεις του πίνακα distance… Αλλά στα τελικά αποτελέσματα η διαφορά αυτή είναι αισθητή???

γενικά δεν παίζει ιδιαίτερο ρόλο

Δεν υπάρχει διαφορά. Φαίνεται από την πολυπλοκότητα του αλγορίθμου αυτό

Ο διάδης είναι κάτω;

ναι ρε γαμωτο, και δεν μπορω να παρω μετρησεις!!!  >:( >:( >:(

Για την δέσμευση μνήμης στην GPU μπορούμε είτε να :
α) κάνουμε όλο τον πίνακα διάνυσμα N*N και να χρησιμοποιήσουμε cudaMalloc( )
β) να δουλέψουμε με 2D πίνακα και να χρησιμοποιήσουμε τις cudaMallocPitch( ) ή cudaMalloc3D( )

Καμία πρόταση ? Έχει διαφορά στην επίδοση ? Το δοκίμασε κανείς ?
Thanks :)

Για όποιον θέλει να χρησιμοποιήσει cudaMallocPitch( ) εδώ τα λέει καλά:
http://stackoverflow.com/questions/16119943/how-and-when-should-i-use-pitched-pointer-with-the-cuda-api (http://stackoverflow.com/questions/16119943/how-and-when-should-i-use-pitched-pointer-with-the-cuda-api)

Καλησπέρα φίλε ,  τρέχει ακόμα η προθεσμία;

Well ... Κανονικά όχι αλλά λέω να την στείλω :P ! Στην χειρότερη δν θα γίνει δεκτή αλλά λέω να το ρισκάρω

(εντελως κατοπιν εορτης)

εχω κολλησει σε ενα θεμα και δεν βρισκω κατι σχετικο καπου.

εγραψα παραλληλο κωδικα της πρωτης περιπτωσης (n*n blocks, 1 thread per block, χωρις shared memory), και συμβαινει το εξης παραδοξο. Μεχρι και n=2^7 δουλευει καλα και επαληθευονται τα αποτελεσματα σε σχεση με τον σειριακο.
Απο 2^8 και μετα ομως, απλα μου επιστρεφει πισω τον αρχικο πινακα, τελειωνοντας πολυ γρηγορα

δοκιμασα επισης υλοποιηση 1 block με n*n threads, το οποιο εμφανιζει το ιδιο προβλημα, αλλα για n=2^6 και μετα.

υποψιαζομαι οτι κατι παιζει με τις μεγιστες δυνατες διαστασεις. εχει κανεις καμμια ιδεα?

Δοκίμασες καμία από τις συναρτήσεις της CUDA που σου επιστρέφουν τα errors κατά την εκτέλεση στη GPU;

Η cudaGetLastError() προσωπικά μου είχε σώσει πολύ ώρα που θα σπαταλούσα σε debugging.

θα το δω :)

Μπορείς επίσης να δεις μέσω της cudaMemGetInfo() το διαθέσιμο χώρο στη GPU, για να είσαι σίγουρος ότι δεν πας να κάνεις allocate παραπάνω χώρο από όσο έχεις διαθέσιμο.

Γενικά αυτό μου φαίνεται ότι χρησιμοποιείς παραπάνω threads απο ότι μπορεί να "σηκώσει"
 ταυτόχρονα. Είχα το ίδιο προβλημα και

Με έσωσε και μένα αν το τρέχεις σε windows υπάρχει άλλο ένα μικρό θεματάκι που μου πήρε μερικές μέρες αλλά το βρήκα
Γενικά υπάρχει τρόπος να τα δεις τα μέγιστα threads με το cudastatus νομίζω (στην περίπτωση που δεν τρέχεις στο pc και να αρκεστείς σε ένα googlαρισμα της κάρτας γραφικών )

Έτσι βλέπεις τι σφάλμα έχεις: