Archi GPU et GPGPU

[Møgluglu]

Mais je pense qu'il faut juste une archi un peu couillue en flottant SP...

N'importe quoi.

La réponse c'est : pour l'arithmétique d'intervalles il faut un Itanium, qui permette de changer le mode d'arrondi sans pénalité.
Et surtout pas un Alpha, ni un Cell, et encore moins un GPU, qui ne supportent pas les arrondis dirigés en hardware.
(encore que pour l'Alpha je n'ai pas tout compris, j'ai l'impression qu'il y a plusieurs générations différentes. Je suis preneur d'une explication sur les modes d'arrondis...)

Sinon j'ai refait le test sur Xeon 5148 Woodcrest en x86_64 (enfin sur 1 de ses cores). Les résultats obtenus m'ont fait me souvenir que Boost à été programmé par des gens très scrupuleux, et que par défaut les opérateurs restorent le mode d'arrondi après chaque appel. Donc j'ai retesté en l'autorisant à garder le même mode tout le temps (noté 2)...
J'obtiens (toujours en ns/op) :

op	P4 1	P4 2	Xeon 1	Xeon 2
x+y	7,75	8,34	89,4	-0,59
x*y	24,44	22,05	94,2	4,19
x^2	19,67	7,15	90,6	2,98
x^5	68,54	63,78	109,7	26,82

(oui, le nombre négatif c'est que ça prend moins de temps que la boucle vide...)

Dans mes souvenirs le Willamette et le Northwood avaient du hardware spécial pour permettre de switcher entre 2 modes de configuration FPU.
J'étais persuadé que ça avait été supprimé sur le Prescott, mais vu les résultats il semble que non.

Prescott forever

[Møgluglu]

Je dois reconnaitre que les résultats m'ont surpris.

Il faut garder à l'esprit que c'est de l'arithmétique sur des intervalles. Les opérations ne sont pas les mêmes qu'en arithmétique conventionnelle.
http://fr.wikipedia.org/wiki/Arithm%...es_intervalles

Ou pour les germanophones :
http://de.wikipedia.org/wiki/Intervallarithmetik

Pour le x² sur CPU comme sur GPU, il est (un peu, pour le CPU) plus rapide que le x*y, et ca doit s'expliquer assez bien vu que x²=x*x, et qu'on a donc qu'un load a faire, par rapport à 2 pour x*y.

Non, ça s'explique surtout par le fait qu'on est sûr que quoi qu'il arrive, x est de même signe que x. Du coup on n'a que 3 cas à considérer (a, b >= 0) :
x = [a, b] -> x^2 = [a^2, b^2]
x = [-a, b] -> x^2 = [0, max(a, b)^2]
x = [-a, -b] -> x^2 = [b^2, a^2]

À comparer avec la multiplication qui, dans l'implémentation de Boost, a 4 niveaux de if imbriqués pour distinguer 13 cas (je fais pas comme ça sur GPU).

Par contre, pour x^5
Si tu en es fier, je suppose que c'est plus compliqué que x*x*x*x*x ? (remarque, les résultats en eux-même sont suffisants pour être fier )

C'est presque aussi simple que (x*x)*(x*x)*x, c'est pour ça que ça va vite
Mais je suis très malhonnête dans mes résultats. D'une part, je massacre la précision sur mon GPU, d'autre part ma fonction n'élève que par des constantes, alors que le code contre lequel je compare le fait pour des variables (certes, pour un compilo assez intelligent ça ne devrait pas faire de différence, mais c'est pas le cas de gcc).

Pour x+x, tu as combien ? (ah oui mais le compilo va faire un mul 2, pas un add...)

Non non, moi quand je lui demande de faire une addition il me fait une addition (enfin au moins tant que c'est moi qui ai programmé l'opérateur d'addition...)

Je gagne entre 10 et 12 ps par opération (en moyenne!) par rapport à x+y. Et entre 20 et 30 ps par rapport à 2 * x.

Tu bosses pour une université française ?

Oui, une petite fac dans une ville paumée mais néanmoins ensolleillée. Comment tu as deviné?

Clair que ta 8500, elle fait un peu pitié (et je ne parle pas ton prescott bien sur).

Sauvons la recherche...

[fefe]

Une petite question quand tu mesure la perf de ton test:
tu testes n operations (x+y par ex) differentes, mesure le temps et divise par n,
ou tu fais le meme x+y nfois et divise le temps par n ?

Dans le premier cas tu laisses intervenir la bande passante de la memoire dans ton equation et part du principe que l'utilisateur additionnera des valeurs qui n'avaient pas ete utilisees precedement, et dans le second cas tu surestimes la capacite du cache du CPU a toujours avoir la donnee dispo.

Le premier cas dorne une borne inf de la perf du cpu, le second cas une borne sup de la perf du cpu, pour l'impact sur le GPU je ne suis pas convaincu qu'il soit aussi important que sur le cpu. Suivant les types de calculs en jeux, entre les 2 methodes il y a 1 a 2 ordres de grandeur de perf (vu que tu passes peu de ns pour faire chaque operation, si tu passais des micros s-mili s par op il y aurait peu de diff).

Sinon pour les budgets dans la recherche, ne m'en parle pas, a priori rien n'a change en 10 ans. Justifier l'achat d'une carte 3D en 98 demandait un investissement personnel et une patience sans bornes (par contre en decembre il fallait de depecher de finir le budget de l'annee histoire qu'on ne te le reduise pas l'annee suivante).

[Møgluglu]

Une petite question quand tu mesure la perf de ton test:
tu testes n operations (x+y par ex) differentes, mesure le temps et divise par n,
ou tu fais le meme x+y nfois et divise le temps par n ?

Je fais n opérations différentes sur des valeurs successives que je calcule au vol (mon compteur de boucle en gros). Je rajoute des fausses dépendances pour empêcher le compilo d'optimiser la boucle, mais ça ne devrait pas empêcher l'exécution en pipeline ou le déroulage partiel.

L'autre endroit où je suis malhonnête (encore), c'est que je me place dans un cas où je n'ai pas de problème de misprédiction (sur CPU) ni divergence (sur GPU) de mes branchements.

Dans le premier cas tu laisses intervenir la bande passante de la memoire dans ton equation et part du principe que l'utilisateur additionnera des valeurs qui n'avaient pas ete utilisees precedement, et dans le second cas tu surestimes la capacite du cache du CPU a toujours avoir la donnee dispo.

Ce que je vends, c'est juste une bibliothèque d'opérateurs. Je dois juste faire en sorte qu'ils soient le plus rapides possible. Mes opérateurs ne font pas d'accès mémoire donc je ne prends pas ça en compte dans mes test. C'est au client de ma bibliothèque de concevoir ses algos pour limiter la consommation mémoire, et à la limite ça ne me concerne pas. S'il est limité par la bande passante, mon travail d'optimisation n'aura servi à rien, mais au moins ce n'est pas moi le maillon faible.

Donc effectivement je mesure une borne sup sur la vitesse.

Par contre l'application au raytracing est là pour voir ce que ça donne dans une vraie appli en pratique.

Le premier cas dorne une borne inf de la perf du cpu, le second cas une borne sup de la perf du cpu, pour l'impact sur le GPU je ne suis pas convaincu qu'il soit aussi important que sur le cpu. Suivant les types de calculs en jeux, entre les 2 methodes il y a 1 a 2 ordres de grandeur de perf (vu que tu passes peu de ns pour faire chaque operation, si tu passais des micros s-mili s par op il y aurait peu de diff).

La 8500 a un ratio bande-passante / puissance de calcul assez impressionnant, on peut saturer les unités de calcul très rapidement pour peu que les accès mémoire soient bien réguliers. Sur le CPU c'est effectivement différent.

Sinon pour les budgets dans la recherche, ne m'en parle pas, a priori rien n'a change en 10 ans. Justifier l'achat d'une carte 3D en 98 demandait un investissement personnel et une patience sans bornes (par contre en decembre il fallait de depecher de finir le budget de l'annee histoire qu'on ne te le reduise pas l'annee suivante).

Voilà, c'est exactement ça.
Mais depuis l'ANR la majorité des financements sont affectés à des projets précis. Du coup on en arrive à déployer des trésors d'imagination pour arriver à raccrocher des travaux à un projet...

[fefe]

Je rajoute des fausses dépendances pour empêcher le compilo d'optimiser la boucle, mais ça ne devrait pas empêcher l'exécution en pipeline ou le déroulage partiel.

Comme je l'avais dit sur le topic a Sam sur la mesure de temps, tu peux t'essayer a mettre des rdpm avant et apres une operation et mesurer la latence d'une seule operation dans ta boucle sans mettre de fausse dependances (tu peux mesurer rdpm op de latence connue rdpm pour le soustraire a ta latence). Apres tu peux regarder le min/max/average de ta latence pour voir l'impact de l'ordonnancement.

Casser les dependances defavorise pas mal les CPUs en general etant donne qu'il aurait probablement execute ce code en partie en parallele avec ce qu'il y avait avant/apres. 7ns/op c'est une 15 aine de cycles sur ta machine, si tu forces des dependances qui n'existent pas avant et apres tu perds beaucoup de perf sur un OOO moderne (je ne dirais pas ca si tu avais 200+cycles/op). Cote GPU aucune idee du cout de la chose en perf, mais etant donne des capacite de reordonnancement plus limitees, je dirais beaucoup moins d'impact.

Edit: Mmm je me rends compte que tu compiles ton code x86 et que je te demande de passer en asm pour mesurer les perfs / oublie .

[Foudge]

fefe> T'as un exemple d'utilisation du rdpm (RDPMC j'imagine) ? Une initialisation quelconque à faire ? J'ai essayé pour voir (en passant par un driver en ring0 sinon exception) et j'obtiens toujours 0 (et 0-0 ça fait inévitablement 0)

[fefe]

Oui il faut d'abord l'initialiser pour dire quel performance monitor tu configures, je suppose que c'est ce qui t'arrive. Je te trouve ca.

[fefe]

Tu configures le numero du PM que tu veux tracker (http://download.intel.com/design/Pen...s/24319202.pdf chapitre 15.6), 79H pour cpu_clck_unhalted, puis tu pointes dans ECX sur ce PM (cf http://download.intel.com/design/Pen...s/24319202.pdf appendix A-10) et RDPMC te retourne la valeur du compteur dans EDX:EAX (LSB dans EAX, le reste dans EDX) http://www.intel.com/design/processo...als/253667.pdf page 4-315.

The performance-monitoring counters are supported by four MSRs: the performance event
select MSRs (PerfEvtSel0 and PerfEvtSel1) and the performance counter MSRs (PerfCtr0 and
PerfCtr1). These registers can be read from and written to using the RDMSR and WRMSR
instructions, respectively. They can be accessed using these instructions only when operating at
privilege level 0. The PerfCtr0 and PerfCtr1 MSRs can be read from any privilege level using
the RDPMC (read performance-monitoring counters) instruction.
NOTE
The PerfEvtSel0, PerfEvtSel1, PerfCtr0, and PerfCtr1 MSRs and the events
listed in Table A-1 in Appendix A, Performance-Monitoring Events are
model-specific for P6 family processors. They are not guaranteed to be
available in future Intel Architecture processors.

Dans ton cas tu ferais un WrMSR de 79H dans PerfEvtSel0, 0 dans ECX au debut, et tu devrais pouvoir faire tes RDPMC pendant l'exec de ton prog. Les details sur tous les parametres configurables pour tes PM sont dans la section 15.6

Sinon sur quel CPU as tu teste ? Les index des PM ont beaucoup change entre P6/P4/Core c'est peut etre de la que vient le pb ?

Pour la necessite de passer en Ring 0 c'est que CR4.PCE est a 0 (il faut etre en ring 0 aussi pour le changer mais ca te permet d'utiliser les PM en dehors de ring 0).

[Foudge]

Merci pour toutes ces infos (même si c'est un peu HS). J'ai testé sur un C2Q avec le PM 0 et 1, ça donnait le même résultat.
Je vais voir comment changer le CR4.PCE car ça pourrait me faciliter les choses.

[Møgluglu]

(Titre du topic réajusté en fonction du contenu)

Avec la fin du NDA sur le RV770, quelques articles intéressants.
HFR
The Tech Report
PC-Watch (Hiroshige Goto)

L'article de Goto a un joli tableau récapitulatif.

En ce qui concerne le GPGPU, on a en hardware un monstre de 1 TFlop pour une taille de die et conso raisonnable. Par contre (comme souvent chez ATI?) le software ne suit pas trop.
AMD a repris le projet Brook de Stanford à son compte, mais en est toujours à débugger le code (il en a besoin) et n'a quasiment pas ajouté de features. Surtout, Brook s'appuie toujours principalement sur la couche graphique, avec HLSL comme langage intermédiaire.
Donc en gros, pour le moment Brook = HLSL - features graphiques , et on se prend encore les limitations "graphiques" dans les dents.

En dessous ils ont CAL, avec un langage intermédiaire niveau assembleur, qui donne accès à un sous-ensemble de fonctionnalités un peu plus grand, mais largement en dessous de ce que le hardware peut faire (mémoires locales, accès direct en RAM sans passer par le cache, features graphiques, etc.)
Donc pour le moment je ne suis pas trop convaincu, à voir avec la prochaine release en juillet/août...

Edit: et contrairement à ce qu'AMD prétend, le RV670 a bien un FMA en double comme le GT200.

[fefe]

J'ai entendu des bruits de couloir disant qu'AMD abandonnait progressivement leur solution actuelle pour OpenCL.
Que penses tu d'OpenCL ?

Sinon une petite question si tu as programme sur des GPUs avant et apres le passage aux scalaires:
Est ce que d'avoir des unites "scalaires" aide particulierement le programmeur plutot que des unites SIMD ? Sachant que au final la contrainte d'executer la meme instruction en parallele sur chaque unite scalaire reste (mais pas les conditions d'alignement des donnees). Je suppose que si tu as un gather et un scatter performants il n'y a pas de difference fondamentale, mais je demande confirmation (je suppose que dans des programmes avec beaucoup de branchements ca peut avoir un interet mais je ne suis pas convaincu que ca change vraiment la vie).

[Møgluglu]

J'ai entendu des bruits de couloir disant qu'AMD abandonnait progressivement leur solution actuelle pour OpenCL.
Que penses tu d'OpenCL ?

Rien, vu que je n'ai rien vu de concret. Ça semble une initiative louable, mais en partant de zéro et fait à la OpenGL (longues discussions avant d'intégrer la moindre feature) ça va certainement prendre des années. Et pendant ce temps les développeurs vont s'habituer à Cuda.
Pour moi il est encore trop tôt pour définir un standard, on n'a pas le recul nécessaire pour faire les bons choix. Le hardware évolue encore beaucoup trop vite et le standard risque d'être déjà à la rue au moment ou il sera finalisé. Mieux vaut expérimenter sur des langages et bibliothèques "jetables", et standardiser a posteriori une fois que c'est stabilisé.

Sinon une petite question si tu as programme sur des GPUs avant et apres le passage aux scalaires:
Est ce que d'avoir des unites "scalaires" aide particulierement le programmeur plutot que des unites SIMD ? Sachant que au final la contrainte d'executer la meme instruction en parallele sur chaque unite scalaire reste (mais pas les conditions d'alignement des donnees). Je suppose que si tu as un gather et un scatter performants il n'y a pas de difference fondamentale, mais je demande confirmation (je suppose que dans des programmes avec beaucoup de branchements ca peut avoir un interet mais je ne suis pas convaincu que ca change vraiment la vie).

Je n'ai pas beaucoup programmé sur les anciens GPU, mais effectivement il n'y a pas de différence fondamentale. L'approche scalaire ("SIMT") simplifie juste la vie pour les branchements, gathers et scatters. L'approche SIMD explicite à la SSE permet d'échanger les données plus facilement d'un thread à un autre (shuffle, dot prod, hadd, sad...).
Le gros avantage du SIMT est de gérer dynamiquement en hardware les branchements et les accès mémoire et de reconnaître ceux qui sont parallélisables, sans que le programmeur/compilo ait besoin de les marquer explicitement en hardware-friendly ou non. Sans compter ceux qui sont hardware-friendly dans 99% des cas mais pas toujours. C'est un peu comme la question des accès non-alignés sur CPU.

Mais on peut convertir mécaniquement un code d'un mode à l'autre. NVidia vient justement d'annoncer un backend SSE pour Cuda.

Et d'un autre côté le GT200 apporte les instructions de voting, qui permettent de faire un AND ou un OR horizontal (si un des threads du groupe / tous les threads du groupe a/ont la donnée à vrai alors...) Ça va dans le sens d'un retour au SIMD explicite, ou au moins de quelque chose entre les deux. C'est le genre de chose qui me fait dire pour la standardisation : wait and see.

[fefe]

Je n'ai pas beaucoup programmé sur les anciens GPU, mais effectivement il n'y a pas de différence fondamentale. L'approche scalaire ("SIMT") simplifie juste la vie pour les branchements, gathers et scatters. L'approche SIMD explicite à la SSE permet d'échanger les données plus facilement d'un thread à un autre (shuffle, dot prod, hadd, sad...).
Le gros avantage du SIMT est de gérer dynamiquement en hardware les branchements et les accès mémoire et de reconnaître ceux qui sont parallélisables, sans que le programmeur/compilo ait besoin de les marquer explicitement en hardware-friendly ou non. Sans compter ceux qui sont hardware-friendly dans 99% des cas mais pas toujours. C'est un peu comme la question des accès non-alignés sur CPU.

C'est un peu le meme probleme que VLIW vs Superscalaire OOO: vaut il mieux extraire le parallelisme statiquement ou dynamiquement, sachant qu'on en trouvera toujours plus dynamiquement, mais que l'avantage est generalement lie au profil des applis.
Si il y a effectivement beaucoup de branchement data dependants et que l'enthropie de ceux-ci est tres elevee, alors effectivement l'approche statique est moins efficace, mais mon impression etait que pour les codes FP portes sur GPGPU n'exhibaient pas trop ce genre de caracteristiques, car ce seront probablement les dernieres applications a vouloir rester sur un CPU generaliste avec une prediction avancee et des schedulers de grande taille.

Donc en fait je ne comprends pas vraiment. Je suis convaincu que les ingenieurs d'ati et nvidia savaient parfaitement ce qu'il faisaient quand ils sont passes au SIMT au lieu de SIMD. Mais pour moi ce n'est pas completement intuitif. Les schedulers qu'ils ont du implementer en hard pour passer a du SIMT ont un cout non negligeable, et le silicone aurait pu etre employe pour mettre des unites de calcul a la place.
Dans un CPU desktop ce choix est assez evident, le gain en performance moyenne est superieur a ce que la surface occuppee par les schedulers et les predicteurs pourrait apporter si transformee en unites d'execution. Le fait que ca soit le cas aussi pour les GPUs me surprend.

[Møgluglu]

Pour moi le SIMT recouvre 2 aspects :
- Le SIMD implicite, ou le fait que le programmeur ne voit qu'un seul thread pour les branchements et les accès mémoire.
Ça on l'avait déjà dans les anciens GPU, en plus d'un SIMD explicite de largeur 4. En SSE, on n'a que le SIMD explicite. C'est à ça que je pensais en écrivant le post précédent. Vu qu'il s'agit juste de grouper ou non les threads en paquets de 4, ça reste du SIMD et ça ne change pas trop la complexité du matériel... Et le compilateur peut utiliser l'un ou l'autre avec des perfs du même ordre.

- L'utilisation de SMT pour masquer la latence, là où les anciens GPU n'utilisaient que du SIMD (implicite). Là effectivement passer du SIMD au SMT oblige à avoir un scheduler en plus.

Mais ils n'avaient pas vraiment le choix, vu que c'est la seule façon de gérer les branchements pas trop inefficacement. Le compilateur ne peut rien y faire.

Sauf si on regroupe les threads dynamiquement en fonction de leur comportement : il y a quelques papiers de recherche là-dessus, mais ça revient en gros à déplacer la complexité du scheduler vers le groupeur de threads et le déplaceur de contextes.
Et en graphisme, non seulement il y a des branchements dans le code des shaders, mais en plus les polygones ont des bords et ne remplissent pas toujours des carrés de 32x32 pixels (même si le Setup Engine sait allouer les threads plus intelligemment que ça, mais ça revient aussi à déplacer la complexité du scheduler).

Et puis la largeur SIMD restante est certainement calculée de façon à ce que le scheduler ne soit pas sur le chemin critique.

[DJ_DaMS]

Et si AMD adoptait Cuda :
http://www.presence-pc.com/actualite...ysX-ATI-30144/
Sans compter que dans l'article de HFR, Damien explique que les caches que ATI a ajouté à son RV770 permettrait l'utilisation de Cuda.

[Møgluglu]

Et si AMD adoptait Cuda :
http://www.presence-pc.com/actualite...ysX-ATI-30144/
Sans compter que dans l'article de HFR, Damien explique que les caches que ATI a ajouté à son RV770 permettrait l'utilisation de Cuda.

Mouais... Je dirais peu probable, Not Invented Here.
Et portage de PhysX n'implique pas portage de Cuda, il suffit juste d'écrire un backend.

Après, techniquement c'est faisable.
Mais est-ce aussi dans l'intérêt de NVidia si jamais il s'avère que les applis Cuda vont plus vite sur ATI?

[fefe]

Oui, si CUDA devient le x86 des GPUs

[Møgluglu]

À propos du scheduler du G80, la seule source officielle sérieuse que j'ai trouvé [1] dit :

Implementing zero-overhead warp scheduling for a dynamic mix of different warp programs and program types was a challenging design problem. A scoreboard qualifies each warp for issue each cycle.
The instruction scheduler prioritizes all ready warps and selects the one with highest priority for issue. Prioritization considers warp type, instruction type, and ‘‘fairness’’ to all warps executing in the SM.

Sur les tests que j'ai fait, j'ai parfois des résultats non déterministes, donc soit leur algo de scheduling est randomisé soit il considère des paramètres que je ne maîtrise pas.

Et en réalité ce n'est pas toujours du "zero-overhead" (au moins sur le G86, faut que je teste sur le GT200). Si je fais tourner un seul warp, il faut attendre 8 cycles entre deux instructions même non-dépendantes au lieu de 4.
Si je fais tourner 8 warps et que j'en bloque 6 sur des accès mémoire, les 2 warps restants n'arrivent qu'à exécuter chacun une instruction tous les 16 cycles, alors que quand ils sont tous seuls ils arrivent à 1 instruction en 8 cycles.

Je n'ai pas réussi à déterminer s'il y avait certains threads qui pouvaient plus facilement être schedulés à la suite l'un de l'autre, apparemment c'est plus compliqué que ça.

Donc oui, le coût du scheduler n'est pas négligeable, ce qui aurait obligé NV à implanter des algos moins chers et moins efficaces que le brute-force.

[1] Lindholm, Nickolls, Oberman, Montrym. NVidia Tesla: A Unified Graphics and Computing Architecture. IEEE Micro Hot Chips 19, 2008

[Alexko]

Tiens, et que penses-tu de ça ?

http://www.ngohq.com/news/14219-phys...hd-3850-a.html

[Møgluglu]

Tiens, et que penses-tu de ça ?
http://www.ngohq.com/news/14219-phys...hd-3850-a.html

Si c'est vrai, ça montre juste que PhysX tourne sur Radeon, pas Cuda.

Après je serai le premier à être content si j'ai tord. Cuda est une très bonne API et le R770 est un très bon GPU...

[Foudge]

Et le driver PhysX pour Radeon il sortirait d'où ? Le gars l'aurait codé tout seul dans son coin, en quelques jours ? Ca pue le fake.

[Alice]

Cuda est une très bonne API

Mouais... en fait je n'ai pas aimé l'approche de NVidia qui balancait la spec de CUDA sans donner explicitement son "modèle" hardware. A l'époque ça faisait plus "je te montre comment l'utiliser mais pas comment c'est fait". Il a fallu attendre 2 ans après la sortie du G80 et des premières versions de CUDA pour voir des articles (dont celui que tu sites il me semble) qui détaille explicitement l'architecture et le pourquoi des choix des modèles de prog sous jacent. Avant pour comprendre un peu comment tout ça était fait il fallait faire bcp de spéculation pas mal de test etc. (et lire Beyond3D :-D). Bref CUDA est une bonne approche mais ce n'est le st graal.

Dans tous les cas l'évangélisation massive de NVidia est insupportable. "CUDA c'est l'avenir ça fait tout"... etc sauf qu'en fait non et il semblerait que bcp l'oubli. Ca peux accélérer rapidement et très efficacement des traitements qui s'y prètent bien. Par contre dès qu'on sort des sentiers battus c'est tout aussi complexe à coder/tweaker qu'un CPU (qui lui à l'avantage d'avoir une viabilité de code plus "prévisible") et ce n'est pas pour autant plus performant. Manipuler des données énormes est par exemple d'un monstreux rare alors qu'un CPU s'en sort bien plus simplement. Bref CUDA et plus généralement le GPGPU à un marché mais NVidia le rend bien plus flambant que ce qu'il est réellement.

En allant plus loin, NVidia a énormément misé sur le GPGPU. 2 ans après leur dernière archie qui révolutionnait le marché, ils sortent un GPU énorme avec des ALU 64bit par si des dénormaux par là...etc tout ça pour faire un malheureux +50% dans les jeux par rapport à l'ancienne génération. Simplement décevant qt on sait que s'est justement leur premier marché. Il me semble que ce n'est qu'une fois le RV770 sorti que les sites spécialisés ont commencé à voir a travers ce brouillard marketing: le GTX280 ne tient pas ses promesses. NVidia a fait un monstre calibré pour le GPGPU, AMD/ATI a juste fait un GPU. Le GTX280 est un bon GPU mais j'en attendais bien plus après 2 années de... rien.

[Møgluglu]

Mouais... en fait je n'ai pas aimé l'approche de NVidia qui balancé la spec de CUDA sans donner explicitement son "modèle" hardware. A l'époque ça faisait plus "je te montre comment l'utiliser mais pas comment c'est fait".
Il a fallu attendre 2 ans après la sortie du G80 et des premières versions de CUDA pour voir des articles (dont celui que tu sites il me semble) qui détaille explicitement l'architecture et le pourquoi des choix des modèles de prog sous jacent. Avant pour comprendre un peu comment tout ça était fait il fallait faire bcp de spéculation pas mal de test etc. (et lire Beyond3D :-D).

Tout à fait d'accord sur ce point. Lecture de Beyond3D et de brevets qui trainent, spéculation et tests, je connais bien
Mais ça c'est la politique de NVidia, ça ne présage rien sur la qualité intrinsèque de Cuda.
(Mais c'est sûr que si on ne sait pas l'utiliser correctement c'est dommage...)

Je maintiens que Cuda est une très bonne API. Le problème, c'est qu'en attendant qu'AMD ait débuggé Brook et que OpenCL soit autre chose que 6 lettres sur du papier, il n'existe aucune vraie alternative.

Dans tous les cas l'évangélisation massive de NVidia est insupportable.

On est d'accord sur l'arrogance de NVidia. Oui, le GPU Computing ne fait pas le café, mais je n'ai pas l'impression que les utilisateurs soient dupes. Utilisateurs qui sont toujours principalement des académiques, et quelques boîtes qui regardent ça de loin et mettent 1 ou 2 stagiaires dessus. Il y en a bien qui ont acheté des Tesla (), mais pas qui ont fait d'investissement aussi massif que NVidia aimerait le faire croire.

En allant plus loin, NVidia a énormément misé sur le GPGPU. 2 ans après leur dernière archie qui révolutionnait le marché, ils sortent un GPU énorme avec des ALU 64bit par si des dénormaux par là...etc tout ça pour faire un malheureux +50% dans les jeux par rapport à l'ancienne génération.

Si on regarde simplement le coût en silicium, ce n'est certainement pas la faute à leur pauvre unité FMA double pour 8 MAD+8 MUL en simple (ce qui par ailleurs est amplement suffisant comme ratio), ni aux dénormaux qui rajoutent juste un peu de latence (un peu... 54 cycles pour un FMA en double quand-même ) si on récupère la circuiterie de conversion int<->double pour les traiter.

Si on considère la répartition de l'investissement en développement, peut-être...
Mais c'est pas moi qui vais me plaindre, les cartes qu'ils arrivent pas à vendre elles finissent chez nous

[Møgluglu]

D'ailleurs, pas sûr que le GT200 soit si bon que ça même en GPGPU (contrairement à ce que j'ai dit la semaine dernière ). Ouais il a la double, les dénormaux... le plus intéressant est peut-être le doublement des registres, mais c'était de toute manière nécessaire pour pouvoir augmenter le nombre de threads pour masquer la latence mémoire qui augmente (417ns sur mon ES contre 325ns sur une 9800GX2).

Au niveau micro-archi l'IPC est strictement identique au G80 à drivers identiques (les fameux MUL sont aussi exploitables sur le G80 que le GT200, c'est juste le compilo et les drivers qui ont éventuellement changé), et la fréquence a baissé. Au final la puissance crête en simple de la GTX 280 est juste 33% supérieure à celle de la 9800GTX+, pour une puce 150% plus grosse.

Reste la bande passante mémoire monstrueuse du bus 512 bits qui est un bon atout sur le calcul scientifique.

En face AMD promet des téraflops en simple, a aussi le FMA en double avec pas mal de gigaflops, et une bande passante pas ridicule sur un GPU assez économe pour en mettre 2 ou 3 (ou 4...) par carte...

Le seul avantage (mais de taille) pour le GPGPU de NVidia, c'est Cuda. Du coup si on avait Cuda pour R600-R700 la situation serait... intéressante
NVidia pourrait tout de même s'en sortir avec une implémentation à 2 vitesses, comme pour Cg.

Après, peut-être que le GT200 va se révéler meilleur en pratique que sur le papier (je n'ai fais que des micro-benchmarks), mais ça serait étonnant... (au mieux sur les applis limités mémoire on retrouvera le ratio des débits mémoires.)

[fefe]

Tu as une liste de latence memoire par GPU en GPGPU ? C'est le genre de truc qui m'interresse... 325-417ns je suppose entre le moment ou tu schedule un acces memoire et le moment ou la donnee est consommable par une operation arithmetique ?

Ca fait un peu drole par rapport aux CPUs avec leur 45-70ns, si tu prends des CPUs avec des acces memoire sur256bits tu montes dans les 100-110ns (a cause du multi-socket).

Je vais essayer de decortiquer un peu un floorplan de GT200 pour voir la quantite de surface qui a ete investie juste pour GPGPU. Sur les G80 en ne comptant que la memoire partagee ajoutee sur chaque SM, on arrivait grosso modo a 1.5-2% du chip (et a ma connaissance ca n'etait utilise que pour GPGPU). Le reste est plus dur a evaluer car a peut aussi avoir ete fait pour ameliorer les perfs du GPU comme le passage a SIMT ou l'augmentation de la taille du banc de registre (plus rapide que l'augmentation du nombre de threads). Sur le GT200 c'est plus facile il y a le double precision et l'augmentation de la taille des multipliers classiques pour ameliorer le support FP, ainsi que le doublement des memoires partagees.

Il faut juste que je trouve une photo de die avec une resolution suffisante...

Edit:
trouve: j'arrive a 2.5% du die pour les FP64. J'aimerais bien voir si ils ont publie des chiffres officiels pour comparer. Les memoires partagees m'ont l'air d'etre plus pres de 1% du die sur le GT200 que de 1.5%.

[Møgluglu]

Tu as une liste de latence memoire par GPU en GPGPU ? C'est le genre de truc qui m'interresse... 325-417ns je suppose entre le moment ou tu schedule un acces memoire et le moment ou la donnee est consommable par une operation arithmetique ?

Je fais le test sur G86 et G80 dés que je peux.
Je fais tourner le code suivant :

bar.sync.u32 0
mov.b32 $r7, %clock
mov.u32 $r1, g[$r5]
mov.b32 $r1, $r1
mov.b32 $r6, %clock

Je force une consommation de la donnée avec le mov r1, r1, sinon le hardware est assez intelligent pour continuer out-of-order...
La mesure est faite sur un warp, avec des accès mémoire "coalescés" (alignés et contigus comme il faut), donc requête de 1024 bits. Si je lis moins de données que ça il prend 8 cycles de plus pour masquer les résultats...

(il faut décerner une médaille au type qui a écrit Decuda)

Ca fait un peu drole par rapport aux CPUs avec leur 45-70ns, si tu prends des CPUs avec des acces memoire sur256bits tu montes dans les 100-110ns (a cause du multi-socket).

C'est sûr que c'est pas les mêmes compromis qui sont fait.
Leurs ALU à ~20 cycles et FPU à ~60 cycles sont pas mal non plus...

Je vais essayer de decortiquer un peu un floorplan de GT200 pour voir la quantite de surface qui a ete investie juste pour GPGPU. Sur les G80 en ne comptant que la memoire partagee ajoutee sur chaque SM, on arrivait grosso modo a 1.5-2% du chip (et a ma connaissance ca n'etait utilise que pour GPGPU). Le reste est plus dur a evaluer car a peut aussi avoir ete fait pour ameliorer les perfs du GPU comme le passage a SIMT ou l'augmentation de la taille du banc de registre (plus rapide que l'augmentation du nombre de threads).

Je ne sais pas bien comment les threads sont créés en rendering, mais je ne pense pas que le nombre de threads soit limité par autre chose que le nombre de registres.

Sur le G80, on peut scheduler 24 warps par SM, avec 256x32 registres par SM ça fait 10 registres par thread, et en pratique plutôt 8 parce qu'on ne peut pas partitionner le banc de registre n'importe comment. 8 c'est très peu quand on n'a pas de pile ni de scratchpad...
Sur le GT200 on arrive à 16 registres/thread, c'est un peu plus confortable mais pas monstrueux non-plus.

Ça c'est si on suppose qu'en graphique on ne switche pas entre plusieurs blocks à la Cuda, sinon ça fait d'autant moins de registres.

Sur le GT200 c'est plus facile il y a le double precision et l'augmentation de la taille des multipliers classiques pour ameliorer le support FP, ainsi que le doublement des memoires partagees.

Cuda continue de me casser mes mul32x32 entiers en mul24x24 et la doc n'a pas changé, donc a priori pas de changement dans les multiplieurs classiques. Et seuls les bancs de registres ont doublé et pas la mémoire partagée.

trouve: j'arrive a 2.5% du die pour les FP64. J'aimerais bien voir si ils ont publie des chiffres officiels pour comparer. Les memoires partagees m'ont l'air d'etre plus pres de 1% du die sur le GT200 que de 1.5%.

Intéressant. Pour ma culture perso, tu places où les FP64? C'est toute la rangée centrale des SM ou juste une partie?
S'agissant de NVidia, faut pas trop espérer de chiffres officiels...

[fefe]

Je fais le test sur G86 et G80 dés que je peux.
Je fais tourner le code suivant :

bar.sync.u32 0
mov.b32 $r7, %clock
mov.u32 $r1, g[$r5]
mov.b32 $r1, $r1
mov.b32 $r6, %clock

Je force une consommation de la donnée avec le mov r1, r1, sinon le hardware est assez intelligent pour continuer out-of-order...
La mesure est faite sur un warp, avec des accès mémoire "coalescés" (alignés et contigus comme il faut), donc requête de 1024 bits. Si je lis moins de données que ça il prend 8 cycles de plus pour masquer les résultats...

(il faut décerner une médaille au type qui a écrit Decuda)

Merci

C'est sûr que c'est pas les mêmes compromis qui sont fait.
Leurs ALU à ~20 cycles et FPU à ~60 cycles sont pas mal non plus...

J'ai des doutes que ca soit les latences de leurs ALU entre l'entree et la sortie de l'ALU (il n'y a pas d'interet fondamental a pipeliner autant), donc je pense qu'une partie non negligeable des 20 cycles vient du scheduler et de l'absence de reseaux de bypass. En tout cas ca me semble raisonable car dans un cpu le hardware necessaire a pouvoir exploiter une ALU avec des latences de 1 clock back to back est nettement plus gros que l'ALU elle meme. En effet le Scheduler resout les dependances et schedule tout ca quelques cycles en avance sur un CPU et recouvrir ces quelques cycles est particulierement complexe (cela force du hardware tres complexe qui resout les dependances en 1 cycle, et vu que ca prend plus d'1 cycle une fois implemente il faut employer un algorithme pipelinable qui est encore plus complexe). Bref sur les 20, j'en vois pas vraiment plus de 5 a faire des calculs... Sur le FP par contre, pas de problemes, surtout quand on traite les denormaux a pleine bande passante ca signifie que c'est dans le pipeline pour tout le monde. Pour partager au mieux le hard ils peuvent pipeliner leur mul jusqu'a avoir un etage de radix par clock (Radix 8, 3 bits / clock, 8 clocks) et probablement 5 etages pour l'addition derriere (je ne vois pas comment pipeliner plus), 4 cycles de pre processing et 4 de normalization et arrondi (et je suis vraiment large, je pipeline vraiment tout) ca donne un multiplieur en 21 clocks et je ne vois pas le moindre interret de pipeliner le hard plus (en pratique ce que j'ai decrit c'est la limite du pipelining, ce n'est probablement pas interressant de le pipeliner au dela de 8-10 clocks vu qu'ils ont un peu plus de 22 FO4/cycle - 10 a 20% de plus que K10/C2D). Donc les 40 clocks qui restent sont probablement la pour masquer d'autres operations comme le scheduling, l'acces aux registres, les denormaux, des masques ?

Je ne sais pas bien comment les threads sont créés en rendering, mais je ne pense pas que le nombre de threads soit limité par autre chose que le nombre de registres.


Sur le G80, on peut scheduler 24 warps par SM, avec 256x32 registres par SM ça fait 10 registres par thread, et en pratique plutôt 8 parce qu'on ne peut pas partitionner le banc de registre n'importe comment. 8 c'est très peu quand on n'a pas de pile ni de scratchpad...
Sur le GT200 on arrive à 16 registres/thread, c'est un peu plus confortable mais pas monstrueux non-plus.

J'etais arrive aux memes calculs de registres / thread. Si tu as une operation de reduction a faire et que tu veux la masquer il te faut 1 registre destination par cycle de latence effective (je compte 1 cycle = frequence a laquelle ce thread schedule) de l'operation de reduction. Donc avec 8 registres on ne recouvre meme pas une reduction en flottant vu qu'il faut quelques registres en plus que 8 pour stocker les sources. 16 m'a l'air suffisant par contre.

Cuda continue de me casser mes mul32x32 entiers en mul24x24 et la doc n'a pas changé, donc a priori pas de changement dans les multiplieurs classiques. Et seuls les bancs de registres ont doublé et pas la mémoire partagée.

Et le traitement des denormaux ? Pour la memoire aprtagee je me suis emmele els pinceaux avec ati qui vient d'en ajouter une. Les bancs de registres me semblent necessaires pour le GPU classique aussi.

Intéressant. Pour ma culture perso, tu places où les FP64? C'est toute la rangée centrale des SM ou juste une partie?
S'agissant de NVidia, faut pas trop espérer de chiffres officiels...

Partie centraledes SM, oui. Pour les dimensions, un peu plus haut qu'un FP24, et la largeur de 2FP24 + leurs registres). On arrive a peu pres a distinguer le bloc en passant une photo de die sous photoshop avec un bon unsharp mask. J'ai compte les pixels et compare au nombre de pixels pour le die complet (je sais c'est pas super scientifique mais bon .

[fefe]

Edit: pour mes calcul de pipelines je pars du principe que le hardware est de la largeur des donnees a calculer et qu'ils ne les cyclent pas plusieurs fois dans du hardware "etroit". Si c'est le cas alors effectivement de telles latences semblent OK.

[newbie06]

Edit: pour mes calcul de pipelines je pars du principe que le hardware est de la largeur des donnees a calculer et qu'ils ne les cyclent pas plusieurs fois dans du hardware "etroit". Si c'est le cas alors effectivement de telles latences semblent OK.

(nOOb alert) Mais s'ils cyclent plusieurs fois, alors le pipeline aurait de nombreux stalls, ou ce serait masque par les schedulers ?

[fefe]

(nOOb alert) Mais s'ils cyclent plusieurs fois, alors le pipeline aurait de nombreux stalls, ou ce serait masque par les schedulers ?

Je pars du principe qu'ils masqueraient en schedulant suffisament de thread. Je ne sais plus ou j'ai vu que leurs unites etaient "double pumped" mais c'est ca qui m'a mis la puce a l'oreille et le fait que je n'arrive pas a comprendre l'interet de pipeliner autant un multiplieur flottant classique. Au final je ne fais que supposer, et vu que je ne passerai pas de temps a essayer de le verif sur du hard existant ca restera des suppositions .