Crunchez vos adresses URL
|
Rejoignez notre discord
|
Hébergez vos photos
Page 5 sur 22 PremièrePremière 1234567891011121315 ... DernièreDernière
Affichage des résultats 121 à 150 sur 658

Discussion: Archi GPU et GPGPU

  1. #121
    Encore du GT-200 :
    http://www.realworldtech.com/page.cf...WT090808195242

    Cette fois on a beaucoup plus de détails, en particulier sur le chargement des instructions (p.7 et suivantes).

    Il explique bien comment on peut avoir un core superscalaire avec un IPC <1

    Ça correspond bien à ce que j'obtiens avec mes tests, à part pour le nombre de ports sur les bancs de registres (bas de la p.9). 16 ports pour chaque banc de 8K pour un total de 128 ports, ça fait beaucoup .
    (Et c'est de là que vient toute l'histoire du MUL manquant : la bande passante des bancs de registres est le bottleneck.)

    En le relisant il a l'air de vouloir dire que c'est la somme des ports des bancs de registres et de la mémoire partagée qui fait dans les 128, mais c'est pas si clair... Et les opérandes peuvent venir aussi du cache de constantes et d'immédiat dans l'instruction. Et on peut partager des ports si 2 opérandes référencent le même registre.

  2. #122
    Mmm le decompte des registres est interressant mais je n'arrive pas aux memes conclusions avec leurs dessins. Si on suppose autant de banks que de SP, cela fait 5 read + 2 write / 2 fast clock (vu que le mul et le MAD sont depipelines).
    Bien entendu il faut des ports "memoire" pour amener les operandes dans les registres a un moment ou a un autre, le minimum etant probablement de 2R+1W par bank (mais c'est une supposition basee sur des connaissances CPU et pas GPU). De cette maniere j'arrive a 7R/3W /bank.
    Si ils n'ont pas 1 bank pas SP je ne comprends pas les raisons de leur choix.

    Pour comparer: Si on prend la RF integer d'un P4 willamette 2 unites au double de la frequence (2R+1W chacune) + 1 load (2R) + 1 store (2R+1W) a la meme frequence on arrive a 12R+5W (a 3GHz en 130nm...).

    Bien entendu on ne parle pas de la meme taille (16x plus gros pour le GPU).

    Au final j'ai l'impression qu'il fait une erreur d'un facteur 2 dans ses calculs, mais je ne comprends pas suffisament bien pour en etre sur. Il montre d'abord que les mul et MAD sont depipelines donc qu au max on peut en avoir que 1 de chaque par 2 cycles, et apres il somme les registres necessaires pour les 2 et multiplie par 2 a cause de la difference de frequence.

    Bref quand je fais mes calculs je ne trouve pas un nombre de ports impressionants a leur register file compare a l'essentiel des microprocesseurs superscalaires de ces ~15 dernieres annees (8R/4W est typique des machines qui peuvent dispatch 4 instructions simultanees, cf un A21164). Par contre ils ont beaucoup de registres.

    Sinon le cout de ports en ecriture est nettement plus elevee que les ports en lecture donc les ajouter ensemble est un peu optimiste si c'est pour mesurer la complexite du circuit.

    PS: sinon l'article est interressant
    fefe - Dillon Y'Bon

  3. #123
    Citation Envoyé par fefe Voir le message
    Bien entendu il faut des ports "memoire" pour amener les operandes dans les registres a un moment ou a un autre, le minimum etant probablement de 2R+1W par bank (mais c'est une supposition basee sur des connaissances CPU et pas GPU).
    Sur GPU c'est pas sûr. Les unités mémoire/front-end de texture n'apparaissent pas sur les schémas, mais a priori il n'y a que 2 ports d'exec, donc on ne peut avoir que 2 instructions en cours d'exécution à un cycle donné (ALU, FPU, load ou store), et chacune peut écrire au max dans 1 destination, donc 2W devrait suffire.

    Au final j'ai l'impression qu'il fait une erreur d'un facteur 2 dans ses calculs, mais je ne comprends pas suffisament bien pour en etre sur.
    Je pense que c'est parce qu'il considère que les ports des bank font 32-bit.
    Si on compte en débit en régime permanent MAD+MUL pendant 4 cycles rapides=2 cycles lents, on a besoin de :
    - MAD: (3R+1W) x 4 cycles x 8 SP x 32 bits
    - MUL: (2R+1W) x 4 cycles x 8 SP x 32 bits

    En supposant 1 bank/SP, le RF doit fournir la bande passante :
    (X*R + Y*W) x 2 cycles x 8 banks x Z bits

    Si Z=32, on tombe bien sur (10R + 4W)x8 banks.

    Mais Z=64 est nettement plus crédible. On ne va pas multiplier par 2 le nombre de ports alors qu'on sait qu'on va toujours les utiliser par paires pour accéder aux entrées 2k et 2k+1. On retombe sur ton calcul.

    Bref quand je fais mes calculs je ne trouve pas un nombre de ports impressionants a leur register file compare a l'essentiel des microprocesseurs superscalaires de ces ~15 dernieres annees (8R/4W est typique des machines qui peuvent dispatch 4 instructions simultanees, cf un A21164). Par contre ils ont beaucoup de registres.
    Le G80/GT-200 a un IPC max de 1 par rapport à l'horloge lente (au moins avec CUDA). Donc on pourrait se contenter de 3R/1W ports de 128-bit. Mais les vertex shaders utilisent des warps de 16 et ont besoin d'une granularité inférieure, donc on est plutôt à 4R/2W.

    Il faut que je refasse mes tests sur le GT-200 pour être sûr, mais jusqu'au G92 on ne pouvait pas alterner des MAD et des MUL en lisant tout les opérandes depuis les registres. Il fallait utiliser des constantes comme opérandes ou réutiliser des registres (élévation au carré) pour atteindre le débit max. Ce qui me fait dire qu'il y a plutôt 4R/2W que 5R/2W.

    Bon. Toute cette théorie c'est beau mais mes tests disent :
    - Pour une instruction donnée, les opérandes sont lus séquentiellement. En latence back-to-back :
    MAD R0, R1, R1, R1 prend 20 cycles (rapides)
    MAD R0, R1, R1, R2 prend 22 cycles
    MAD R0, R1, R2, R3 prend 24 cycles

    (Sur le GT-200. Sur G80-G92, il y a 2 cycles de moins)

    - Ce n'est pas pipeliné, le MAD suivant un MAD à 3 source démarre 6 cycles après. Je n'ai pas encore réussi à capturer un MUL/Spécial juste après un MAD3 mais je suppose que l'autre unité a aussi son propre port de lecture, éventuellement son port d'écriture.

    (non, parce qu'écrire un microbenchmark en SPMD pour arriver à forcer le scheduler à envoyer un MUL juste après un MAD entre deux lectures de l'horloge interne sans mixer d'autres instructions, c'est bien prise de tête )

    Avec CUDA (et les pixel shaders) ça se comporte donc comme 2R/1Wx128-bit...
    Ça peut être implémenté avec 4R/2Wx64-bit, mais c'est un peu du gâchis.

    Les vertex shaders pourraient profiter de ports de 64-bit plus nombreux, mais il sont limités par l'issue rate de 1 de toute façon.

  4. #124
    J'ai une nouvelle hypothèse.

    Sur le floorplan du GT-200, on voit 8 paires de blocs pour les bancs de registres. Je suppose que veut dire qu'il contient 8 banks.
    On pense tout de suite 8 banks pour 8 SP, 1 bank pour chaque SP, logique non?

    Sauf que la doc de CUDA dit :
    Generally, accessing a register is zero extra clock cycles per instruction, but delays
    may occur due to register read-after-write dependencies and register memory bank
    conflicts.
    [...]
    The compiler and thread scheduler schedule the instructions as optimally as possible
    to avoid register memory bank conflicts. They achieve best results when the number
    of threads per block is a multiple of 64.
    Pourquoi il me parlent de conflits si chaque SP a sa bank dédiée? Si ça se trouve c'est pas comme ça que ça marche.

    Jusqu'ici, j'ai surtout fait mes tests avec peu de threads (64 ou 128), parce qu'après j'ai plus assez de Crayolas pour dessiner les pipelines .
    Dans ces tests-là, ça se passe comme si j'avais 1 seul port dans le RF.
    Maintenant si je lance plus de threads (256), oh miracle, des nouveaux ports qui apparaissent.
    256 threads, ça fait 8 warps. Tiens, autant que de banks dans le RF...

    Donc la nouvelle hypothèse de travail, c'est que le RF est divisé en 8 banks de 256x512-bit, voire 128x1024-bit. Entre le RF et les SP on a un gros crossbar qui scatter les 512-bit/1024-bit lus en 8 SPx2 clocksx32-bit/64-bit.

    Chaque bank appartient exclusivement à un warp ou groupe de warps (jusqu'à 4). Tiens et le nombre de registres adressables par un thread c'est 128.
    Chaque bank étant en 1R/1W, les opérandes d'une instruction doivent bien êtres lus séquentiellement. Par contre, c'est overlappable avec les lectures d'opérandes d'autres threads, mais à condition d'avoir assez de threads.

    (j'ai raté ma vie, j'aurai du faire biologiste pour étudier les puces, ou physicien au CERN pour faire des expériences... )
    Dernière modification par Møgluglu ; 12/09/2008 à 20h34.

  5. #125
    Citation Envoyé par Møgluglu Voir le message
    (j'ai raté ma vie, j'aurai du faire biologiste pour étudier les puces, ou physicien au CERN pour faire des expériences... )
    Ou "Reverseur" pour Intel ?
    Qui sait, cooptation se pourrait-il ?

  6. #126
    Citation Envoyé par Møgluglu Voir le message
    J'ai une nouvelle hypothèse.

    Sur le floorplan du GT-200, on voit 8 paires de blocs pour les bancs de registres. Je suppose que veut dire qu'il contient 8 banks.
    On pense tout de suite 8 banks pour 8 SP, 1 bank pour chaque SP, logique non?

    Sauf que la doc de CUDA dit :


    Pourquoi il me parlent de conflits si chaque SP a sa bank dédiée? Si ça se trouve c'est pas comme ça que ça marche.

    Jusqu'ici, j'ai surtout fait mes tests avec peu de threads (64 ou 128), parce qu'après j'ai plus assez de Crayolas pour dessiner les pipelines .
    Dans ces tests-là, ça se passe comme si j'avais 1 seul port dans le RF.
    Maintenant si je lance plus de threads (256), oh miracle, des nouveaux ports qui apparaissent.
    256 threads, ça fait 8 warps. Tiens, autant que de banks dans le RF...

    Donc la nouvelle hypothèse de travail, c'est que le RF est divisé en 8 banks de 256x512-bit, voire 128x1024-bit. Entre le RF et les SP on a un gros crossbar qui scatter les 512-bit/1024-bit lus en 8 SPx2 clocksx32-bit/64-bit.

    Chaque bank appartient exclusivement à un warp ou groupe de warps (jusqu'à 4). Tiens et le nombre de registres adressables par un thread c'est 128.
    Chaque bank étant en 1R/1W, les opérandes d'une instruction doivent bien êtres lus séquentiellement. Par contre, c'est overlappable avec les lectures d'opérandes d'autres threads, mais à condition d'avoir assez de threads.
    Quand je parlais d'une Bank par SP je ne parlais pas d'une bank privee/dediee. En general quand on bank une structure memoire on essaye de reduire au maximum les conflits, donc on met au moins autant de banks que d'agents qui peuvent y acceder. Bien entendu tous les agents peuvent acceder a toutes les banks, en general l'addressage est fait a travers une hash fonction (souvent un modulo de base), et on essaye d'optimiser le code et la hash fonction de maniere a ce qu'il soit possible d'avoir tout le monde en train de travailler en parallele sans conflits.
    On peut reduire le nombre de banks par rapport aux agents qui accedent seulement si on sait que statistiquement tout le monde n'accedera pas en parallele, et l'augmenter quand on sait que la pression est tres elevee et qu'il est difficile d'optimiser le code pour garantir peu de conflits.

    Donc ici a mon avis tu as effectivement une bank devant chaque SP, avec un gros reseau de bypass entre chaque bank a latence non uniforme, et une petite hash fonction qui repartit ca aux petits oignons.

    En general dans la litterature ce type de mecanisme est propose pour les caches sous l'appellation NUCA( non uniform cache access), et je suis certain qu'il y a pas mal de litterature ancienne la dessus pour les bancs de registre (pour les caches c'est un sujet d'actualite, pour les registres je m'attends a trouver ca dans des papiers des annees 80 parlant de machines paralleles).

    Pour ce qui est des banks dediees au Warps, c'est une implementation specifique de la hash fonction sus citee qui prend juste le numero du warp pour trouver la bonne bank.

    Cela ne me semble pas etre une bonne idee pour ce qui est de la repartition du load dans le bypass (une fonction nettement plus random te permet d'optimiser la bande passante de ton reseau d'interconnexion), mais c'est une possibilite quand meme.

    Bon je dois y aller, je continuerai a y reflechir plus tard.
    fefe - Dillon Y'Bon

  7. #127
    Citation Envoyé par fefe Voir le message
    Donc ici a mon avis tu as effectivement une bank devant chaque SP, avec un gros reseau de bypass entre chaque bank a latence non uniforme, et une petite hash fonction qui repartit ca aux petits oignons.
    D'accord! Je comprend mieux le dessin de Hiroshige Goto, avec le crossbar au milieu :
    http://pc.watch.impress.co.jp/docs/2.../kaigai446.htm

    Par contre pour la latence non uniforme j'ai des doutes. Au début j'avais fait pas mal de tests en lançant les mêmes instructions en changeant juste les numéros de registres, et je n'ai jamais observé de variation de latence.

    La doc de CUDA semble dire que c'est le scheduler de threads qui s'arrange pour éviter les conflits, ça me semble raisonnable.

    En general dans la litterature ce type de mecanisme est propose pour les caches sous l'appellation NUCA( non uniform cache access), et je suis certain qu'il y a pas mal de litterature ancienne la dessus pour les bancs de registre (pour les caches c'est un sujet d'actualite, pour les registres je m'attends a trouver ca dans des papiers des annees 80 parlant de machines paralleles).
    Merci, je vais faire un peu d'archéologie .

    Pour ce qui est des banks dediees au Warps, c'est une implementation specifique de la hash fonction sus citee qui prend juste le numero du warp pour trouver la bonne bank.

    Cela ne me semble pas etre une bonne idee pour ce qui est de la repartition du load dans le bypass (une fonction nettement plus random te permet d'optimiser la bande passante de ton reseau d'interconnexion), mais c'est une possibilite quand meme.
    Pour une mémoire totalement partagée comme un cache, je comprend qu'il faille randomiser pour répartir le trafic.
    Mais en SIMT les motifs d'accès sont très contraints, et les warps ne sont pas schedulés aléatoirement. Il peut y avoir moyen de faire encore mieux qu'un random parfait en exploitant au contraire les régularités.

    Si on fait entrer les numéros de registres dans la fonction de hashage, les instructions vont avoir une latence variable à la fois en fonction des leurs opérandes et en fonction des autres instructions schedulées en même temps. Ça doit être gérable, mais ça complique le scheduler et le collecteur d'opérandes.

    Si on ne considère que les numéros de warp, le scheduler va avoir plus de facilité à choisir à chaque cycle une instruction qui ne cause pas de conflit (en fait, on a toujours des conflits qui font qu'on doit lire les opérandes séquentiellement, ça fait une latence plus grande mais toujours prédictible). Vu que la politique de scheduling est basée sur du round-robin, on évite déjà naturellement les conflits en évitant de scheduler le même warp 2 fois de suite. Et on peut déplacer toute la gestion des conflits dans le scheduler, en étant sûr que quand une instruction est schedulée elle ne rencontrera pas de conflits.

    Enfin, hypothèse de travail hein...

  8. #128
    Si on prend juste le numero de warp pour determiner la bank, as tu une garantie (ou une forte probabilite) qu'il n'y aura pas 2 acces simultanes? Si oui c'est effectivement un tres bon choix de hash fonction.

    Pour ce type d'acces on ne prend jamais de fonction de hash random, mais une fonction qui a des bonnes proprietes spaciales et temporelles:
    - a un instant donne la plupart des banks sont utilisees (optimiser l'utilisation des banks)
    - les acces vers une bank dans le temps sont repartis uniformement entre les diverses sources (optimiser l'utilisation du reseau d'interconnection)

    Bien entendu cela necessite de comprendre a l'avance les patterns d'acces mais ici c'est tout a fait le cas.

    Pour ce qui est de l'acces non uniforme, cela va dependre de la topologie du reseau d'interconnection. Il est possible de mettre un crossbar qui presentera une latence homogene. C'est plus simple pour le programmeur, mais le design est plus complexe et moins performant: la latence du reseau correspond a la latence du pire des cas au lieu d'etre proche de la moyenne du pire et du meilleur cas si tu as un reseau a latence non uniforme.

    Tu sembles dire qu'ils n'ont pas pris l'option a latence variable, et je comprends que ca aurait probablement rendu le probleme plus difficile pour la generation de code, surtout que ameliorer la latence moyenne n'a pas vraiment d'interet dans leur cas vu que leur architecture est construite autour du nombre de threads approprie pour recouvrir les latences.

    Sinon en regardant a http://pc.watch.impress.co.jp/docs/2...kaigai_11l.gif semble confirmer mes hypotheses.

    En y reflechissant un peu plus, le probleme de scheduling/allocation des registres est assez simple avec une latence uniforme, mais le scheduling avec des lectures/ecritures non uniformes devient un cauchemard. Donc il est quasi certain que les signaux soit latches pour arriver physiquement aux register banks en meme temps et latches au retour pour les lectures pour arriver aux unites d'executions de maniere synchrone.
    Si c'est effectivement un full crossbar ils n'ont pas a se soucier du scheduling de leurs messages sur le reseau d'interconnection. ca coute cher, mais ca simplifie la tache.
    Dernière modification par fefe ; 13/09/2008 à 22h45.
    fefe - Dillon Y'Bon

  9. #129
    Citation Envoyé par fefe Voir le message
    Si on prend juste le numero de warp pour determiner la bank, as tu une garantie (ou une forte probabilite) qu'il n'y aura pas 2 acces simultanes? Si oui c'est effectivement un tres bon choix de hash fonction.
    Vu que l'archi est pensée pour avoir beaucoup de warps en même temps, que quand on en a peu ils rencontrent déjà d'autres types de conflits (apparemment le fetch se fait en round-robin sur 8 warps mini), oui.

    Mais je ne suis pas sûr de ça et ça n'explique pas tout.

    Le problème, c'est qu'en cas de conflit éventuel c'est évité a priori par le scheduler qui va juste choisir une autre instruction. Du coup je vois que telle instruction a été schedulée avant telle autre, mais je ne sais pas si c'est dû à un conflit dans les bank de registres ou à la phase de la Lune.
    Parfois j'ai quand-même des bulles dans mon pipeline, probablement parce que le scheduler ne trouve aucune instruction prête, mais je ne sais pas dire pourquoi.

    Par exemple, un scheduling possible de :
    Code:
    MOV R7, clock
    MUL R3, R7, R7
    MOV R3, clock
    BAR.SYNC  // barrière de synchro


    Chaque couleur correspond à 1 warp. A et B sont les 2 ports d'exec. Sauf gourage, les barres sont de la bonne longueur (lecture + exec + écriture). Les bordures sont autour des instructions de lecture d'horloge dont je suis sûr de la position, les autres sont placées au pif.

    Et pour :
    Code:
    MOV R7, clock
    MAD R3, R1, R2, R0
    MOV R3, clock
    BAR.SYNC


    Le seul truc que j'arrive à dire, c'est qu'il y a probablement 1 seul port d'écriture partagé entre les unités (un bus plutôt qu'un crossbar), vu que j'arrive toujours à trouver des schedulings où je n'ai jamais 2 instructions terminant au même cycle.



  10. #130
    Quand je regarde tes 2 graphs il y a effectivement une tonne de trucs que je ne comprends pas.
    A priori deja il ne semble pas y avoir de classement par age des warps lorsque plusieurs sont elligibles pour un slot. Ensuite il semble qu'il y a 2 schedulers 1 pour A et 1 pour B plutot qu'un unifie (ou alors des priorites en fonctions des ports), sinon je ne vois pas comment le bleu B aurrait reussi a passer devant le rose A au cycle 20 (surtout qu'il interromp une sequence qui jusque la allait tres bien).
    Dans le test du MAD je suis perplexe de voir le rouge passer devant le bleu, et comment le MAD grille la priorite au warp orange du MOV.

    Je pense qu'une des infos qui manquent est la date a laquelle chaque instruction est arrivee dans le scheduler, ca pourrait aider a expliquer certaines des inversions bizarres. A priori ca sent l'algorithme de scheduling qui a un pointeur qui bouge de maniere predeterminee afin de garantir l'absence de deadlock (genre un round robin a la con). De mon experience c'est souvent ce qui justifie les comportements un peu bizarres dans les schedulers, en effet il y a 2 moyens de garantir qu'il n'y aura pas de starvation: tenir compte de l'age de l'instruction dans le scheduling, ca demande une matrice d'anciennete et complexifie pas mal le hard, ou alors avoir une selection des candidats prets par round robin, la on est aussi garanti que personne ne restera bloque mais on a des changements d'ordre un peu bizarre.
    Si je prend ma 2 eme hypothese et ton exemple du MAD, si on suppose que le MAD est arrive dans le scheduler apres que le Warp jaune du mov ait ete lance, on a maintenant tous les warps du MAD et le WARP orange du mov qui sont prets. Pas de bol le pointeur serait passe juste apres le warp orange (si on suppose un round robin de 8, vu que le jaune a rate un slot juste au dessus il est possible que le pointeur l'ait depasse et donc envoie tout les warps du MAD avant de revenir au warp orange du mov).

    C'est un peu tire par les cheveux parce que ca n'explique quand meme pas pourquoi le rouge double le bleu dans le warp du mad mais ca pourrait explique un certain nombre de phenomenes.

    Si on retourne au cas du MUL il n'y a pas le meme type de conflits sur les ports d'exec comme dans le cas du MAD mais on voit clairement que la priorite est donnee a l'envoi d'un warp sur le port oppose.

    PS: d'ou sors tu ces graphes ? Je ne suis aps familier avec les outils dispos pour optimiser sur GPU .
    fefe - Dillon Y'Bon

  11. #131
    Citation Envoyé par fefe Voir le message
    Quand je regarde tes 2 graphs il y a effectivement une tonne de trucs que je ne comprends pas.
    A priori deja il ne semble pas y avoir de classement par age des warps lorsque plusieurs sont elligibles pour un slot. Ensuite il semble qu'il y a 2 schedulers 1 pour A et 1 pour B plutot qu'un unifie (ou alors des priorites en fonctions des ports), sinon je ne vois pas comment le bleu B aurrait reussi a passer devant le rose A au cycle 20 (surtout qu'il interromp une sequence qui jusque la allait tres bien).
    Mmmh... Ça irait à l'encontre du peu d'info venant de NV. Officiellement, on a 1 scheduler qui choisit les instructions suivant un score calculé à partir de l'age du warp, du type d'instruction et des phases lunaires. Si ça se trouve le calcul du score est même microcodé et change à chaque version des drivers.
    Que le port B soit prioritisé par rapport au A n'est pas choquant. Dans un code de calcul le port B est sous-utilisé, donc on a intérêt a l'utiliser dés que c'est possible.

    Si je prend ma 2 eme hypothese et ton exemple du MAD, si on suppose que le MAD est arrive dans le scheduler apres que le Warp jaune du mov ait ete lance, on a maintenant tous les warps du MAD et le WARP orange du mov qui sont prets.
    Tu supposes que les instructions sont fetchées une seule fois pour tous les warps qui l'exécutent? Je n'avais pas pensé à ça, ça pourrait expliquer des choses, comme la limite à 1 nouvelle instruction tous les 8 cycles. Faut que je réfléchisse... (ça m'arrive)

    Sinon pour dessiner mes pipelines, j'ai commencé par placer les instructions timer que je savais placer, puis j'ai mis les autres instructions là où c'était possible. J'ai choisi arbitrairement le scheduling qui était le plus proche du round-robin, mais il y en a plein d'autres possibles qui respectent les contraintes, et le mien est sûrement faux.

    En voilà un dont je suis sûr à 100% (à part pour le bar.sync final, mais c'est histoire de boucher les trous) :


    J'ai toujours du mal à voir la moindre logique/régularité dans l'ordonnancement...

    PS: d'ou sors tu ces graphes ? Je ne suis aps familier avec les outils dispos pour optimiser sur GPU .
    Du pseudo-Excel d'OpenOffice, outil bien connu de tous les optimiseurs sur GPU.
    D'ailleurs je cherche un soft qui me permette de dessiner un pipeline comme ça à partir d'un fichier avec les dates de début/fin/durée, et surtout avec lequel je puisse éditer graphiquement mon pipeline.

    NVidia ont certainement tout plein d'outils de profiling et de debugging en interne, mais bon même si j'étais cambrioleur ils sont un peu loin...
    (et puis à ce moment-là autant voler les specs directement sans se prendre la tête)

  12. #132
    Citation Envoyé par Møgluglu Voir le message
    Mmmh... Ça irait à l'encontre du peu d'info venant de NV. Officiellement, on a 1 scheduler qui choisit les instructions suivant un score calculé à partir de l'age du warp, du type d'instruction et des phases lunaires. Si ça se trouve le calcul du score est même microcodé et change à chaque version des drivers.
    Que le port B soit prioritisé par rapport au A n'est pas choquant. Dans un code de calcul le port B est sous-utilisé, donc on a intérêt a l'utiliser dés que c'est possible.
    Disons que pour pouvoir donner cette priorite a B il faut maintenir l'equivalent de 2 files en interne, meme si le scheduler est unifie il n'est pas homogene (a moins que la priorite soit toujours dans le meme sens en quel cas il n'y a effectivement pas besoin de 2 files).

    Je n'avais pas compris que tu dessinais le scheduling a partir des infos que tu avais (trop habitue a avoir des outils qui me ressortent ce genre de trucs je suppose). Les ordonnancements que tu montrais allaient a l'encontre de la notion d'age dans certains cas (ou c'est limpression que j'en ai eu)
    Tu supposes que les instructions sont fetchées une seule fois pour tous les warps qui l'exécutent? Je n'avais pas pensé à ça, ça pourrait expliquer des choses, comme la limite à 1 nouvelle instruction tous les 8 cycles. Faut que je réfléchisse... (ça m'arrive)
    Je n'ai pas reflechi, ca m'a semble logique, mais je n'ai rien pour le supporter.

    Sinon pour dessiner mes pipelines, j'ai commencé par placer les instructions timer que je savais placer, puis j'ai mis les autres instructions là où c'était possible. J'ai choisi arbitrairement le scheduling qui était le plus proche du round-robin, mais il y en a plein d'autres possibles qui respectent les contraintes, et le mien est sûrement faux.

    En voilà un dont je suis sûr à 100% (à part pour le bar.sync final, mais c'est histoire de boucher les trous) :
    http://tof.canardpc.com/preview/c475...ae9038e581.png

    J'ai toujours du mal à voir la moindre logique/régularité dans l'ordonnancement...



    Du pseudo-Excel d'OpenOffice, outil bien connu de tous les optimiseurs sur GPU.
    D'ailleurs je cherche un soft qui me permette de dessiner un pipeline comme ça à partir d'un fichier avec les dates de début/fin/durée, et surtout avec lequel je puisse éditer graphiquement mon pipeline.

    NVidia ont certainement tout plein d'outils de profiling et de debugging en interne, mais bon même si j'étais cambrioleur ils sont un peu loin...
    (et puis à ce moment-là autant voler les specs directement sans se prendre la tête)
    Oui mais ce genre d'outils sont generalement bases sur la simulation, recuperer ce genre d'infos de hard en train de tourner est assez difficile sans perturber le tout et tres couteux.
    Sinon pour ton truc, c'est assez simple a faire en vba pour excel pour dessiner les pipelines, mais... oui je sors (je suis passe par la pour automatiser des dessins de pipeline donc je parle en connaissance de cause)

    D'ailleurs j'y pense c'est un bon exercice de TP, de faire ecrire un petit scheduler avec quelques contraintes de base, genre nombre de ports dans les registres/ fonction de hash pour eviter les bank conflicts/ aged base/ priorite systematique a une unite. Au final tu peux sortir tres facilement du texte (facon asci art) qui decrit l'exec de ton pipeline. Un warp par ligne, un caractere par cycle . J'en vois qui vont m'aimer d'ici.
    Dernière modification par fefe ; 15/09/2008 à 20h05.
    fefe - Dillon Y'Bon

  13. #133
    C'est normal que j'ai mal à la tête ?
    Il faudrait un forum X86 ultra advanced GT-PRO-XTX
    http://valid.x86-secret.com/cache/banner/313021.png

  14. #134
    Citation Envoyé par DJ_DaMS Voir le message
    C'est normal que j'ai mal à la tête ?
    Oui c'est normal, moi aussi.
    J'y comprends rien à ce !@#$%^ de GPU.

    Sinon bonne idée le sujet de TP, faut que j'arrive à récupèrer les TP d'archi...

  15. #135
    j'ai aussi mal à la tête en vous lisant :D

  16. #136
    Moi j'ai pas encore osé, j'essaierai peut-être demain :D

  17. #137
    Citation Envoyé par dandu Voir le message
    j'ai aussi mal à la tête en vous lisant :D
    Ca doit etre a cause des couleurs flashy dans les dessins de pipelines du GT200 .
    fefe - Dillon Y'Bon

  18. #138
    Citation Envoyé par fefe Voir le message
    Ca doit etre a cause des couleurs flashy dans les dessins de pipelines du GT200 .
    A priori, je ne vois pas d'autre raison

  19. #139
    C'est super instructif. J'ai tout compris.
    Mes propos n'engagent personne, même pas moi.

  20. #140
    Citation Envoyé par Neo_13 Voir le message
    C'est super instructif. J'ai tout compris.
    Cool, explique moi.

    Enfin déjà j'ai compris pourquoi on comprend pas. Mes dessins sont faux. J'ai dis que les opérandes étaient lus séquentiellement, donc pour un MAD l'exécution commence au cycle 3, alors que pour un MOV elle commence au cycle 1. Du coup il faut décaler les MAD de 2 cycles vers le bas.

    Exemple avec un MAD à 3 opérandes et 4 warps :


    Le film :
    La première instruction MOV est fetchée, les warps bleus et rouge démarrent.
    Au cycle 8 le MAD est fetché-décodé, et vert démarre (priorité à l'ancienneté).
    (j'ai des raisons de croire que le front-end délivre 1 instr/8 cycles).
    Violet ne peut pas être schedulé avant le cycle 12. Par contre dés le cycle 10 le MAD bleu peut démarrer, vu qu'il n'utilisera pas les ressources d'exec avant le cycle 16. Du coup violet se fait griller sa place et doit attendre au moins le cycle 18 pour démarrer.

    Au cycle 14, le MAD rouge aimerait démarrer, mais il y a un conflit en lecture dans les registres avec le bleu. Il démarre au cycle 16, et passe devant le pauvre violet.
    Même scénario pour vert qui re-grille la priorité de violet.
    Enfin, au cycle 30, les souffrances de violet s'achèvent par son exécution.

    Après c'est moins clair. Bleu pourrait démarrer au cycle 34, mais doit attendre 36 pour une raison inconnue.
    Rouge démarre au cycle 40 dés que sa dépendance est résolue.
    Violet tente de rattraper son retard en démarrant son MAD. Ce qui retarde vert, qui doit attendre la libération du port A.
    Enfin, violet démarre dés que possible au cycle 66.

    C'est peut-être de la spéculation sauvage comme à la bourse, mais là j'ai presque l'impression que je peux comprendre.
    Dernière modification par Møgluglu ; 16/09/2008 à 18h29.

  21. #141
    Interview de développeurs du RV770 par Anand :
    http://www.anandtech.com/video/showdoc.aspx?i=3469

    Des points intéressants comme la scalabilité, les compromis entre redondance, déclassage et yields, la conso, une allusion aux bancs de registres plus locaux...

  22. #142
    Tres bonne interview effectivement.
    fefe - Dillon Y'Bon

  23. #143
    Très intéressant, ça faisait longtemps que je n'avais pas été aussi happé par un article, au point de le lire d'un coup quand je devrais être en train de dormir.

  24. #144
    C'est prenant comme article, et comme le dit Anand, humanisant quant aux designs d'architecture. Franchement agréable à lire en plus.

  25. #145
    Pas mieux, en effet. Ceci dit, ça reste étrange comme démarche de la part d'ATI, non ?

  26. #146
    Pourquoi ? Ça leur fait un joli coup de pub gratuit...

  27. #147
    Citation Envoyé par Alexko Voir le message
    Pourquoi ? Ça leur fait un joli coup de pub gratuit...
    Demande à Sam tiens, tu verras si c'est pas surprenant de la part d'AMD
    (je dirais que ca l'est moins de la part d'ATI cependant )

  28. #148
    Ah par rapport à leurs habitudes et leur historique de RP toutes pourries ? Oui de ce point de vue, c'est clairement surprenant :D

  29. #149
    Vraiment super comme article, je viens de le lire et c'est le genre d'histoires que j'aimerai voir plus souvent.
    "La vaseline, c'est un truc que j'utilise systématiquement" - vf1000f24

  30. #150
    J'aimerais bien voir ce genre d'article avec comme interviewes des inges qui developpent des CPU (Intel, AMD, ARM, etc.).

Règles de messages

  • Vous ne pouvez pas créer de nouvelles discussions
  • Vous ne pouvez pas envoyer des réponses
  • Vous ne pouvez pas envoyer des pièces jointes
  • Vous ne pouvez pas modifier vos messages
  •