Discussion :

[Rémi Coquet]

@yamashi:

Ok, n'y à t'il cependant pas de "problèmes" l'ors des échanger de gros blocs de données entre CPU et GPU ?

- Comment coopèrent DMA/Kernel streaming/CUDA ?

- Pour l'encodage vidéo, je comprends bien que le fait de décoder "in place" pour un renderer situé dans la carte graphique soit de facto-plus rapide mais qu'en est-il du cas de la "récupération" des données par le CPU ?

- J'entends bien, que continuer à traiter les données dans le GPU serait plus rapide évidement, mais si le CPU gère le réseau et uniquement le réseau, si je désire balancer l'image encoder à d'autres machines via réseau quid de l'écroulement des perfs de µPaul ET de µJacques..?

[yamashi]

Pour ce qui est du transfert CPU <-> GPU sur les dernières carte graphiques, la bande passante de la memoire device est de 140-160GB/s après tout dépend du bus PCIe.
Je ferais un test de la bande passante et de la latence de mes GPU demain.

[Rémi Coquet]

Ok. J'imagine que le "lieu" de la décompression joue aussi:

- Est-il possible d'attribuer une surface mémoire système ?

- Les perfs changent-elles ?

[yamashi]

Il y a belle et bien un système d'émulation.
Sinon je viens de faire un test de bande passante de ma 9600GT et voila ce que j'obtient :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
 
Quick Mode
Host to Device Bandwidth for pageable memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    1639.2
 
Quick Mode
Device to Host Bandwidth for pageable memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    1673.2
 
Quick Mode
Device to Device Bandwidth for pageable memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    43401.5

Et avec de la memoire pinned :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
 
Quick Mode
Host to Device Bandwidth for pinned memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    3195.6
 
Quick Mode
Device to Host Bandwidth for pinned memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    3242.6
 
Quick Mode
Device to Device Bandwidth for pinned memory.
Transfer Size (Bytes)    Bandwidth(MB/s)
33554432                    43383.4

Je testerais d'autres mode voir si il y a de gros changement de perf.
Sinon pour ce qui est des serveurs et des cartes graphiques il est effectivement dur de trouver des hébergeurs qui proposent ce genre de serveur mais si tu vas jusqu'à utiliser CUDA pour un serveur c'est que tu as un gros projet et que tu peux acheter tes serveurs et ensuite prendre un service de housing. Pour les serveurs supportant cuda regarde les Tesla S870 et S1070 ce sont 4 GPU dans un 1U.

- Est-il possible d'attribuer une surface mémoire système ?

Oui.

- Les perfs changent-elles ?

Si tu parles de la pinned memory effectivement, 2 fois plus rapide pour le transfert CPU <-> GPU.

Pinned :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

cutilSafeCall( cudaMallocHost( (void**)&odata, memSize ) );

Pageable :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

data = (unsigned char *)malloc( memSize );

Pour OpenCL effectivement ce n'est pas aussi intéressant que CUDA pour le moment du moins.

[Rémi Coquet]

Il me semblait aussi que

140-160GB/s

c'était un peu optimiste mais c'est déjà très beau.

-Donc pinned mémory, quelles différences/limitations/intérêts (à part la manière de réserver la RAM) entre les deux modes ?

- Le transfert inter-cartes semble donc le plus intéressant, perturbe t'il DMA ou les autres flux PCIe ?

- Pour OpenCL, utilise t'il un logiciel assembleur dédiée ?

- Comment OpenCL peut-il générer du code à la fois pour NVIDIA et AMD ?

- La compression hard (H264) n'est-elle accessible que via CUDA ?

- Permet t'elle de travailler avec une key frame par frame etc. ou est-ce que les différents mode sont "gelés"

(j'arrête là parce que nous y passerions la nuit )

[yamashi]

140-160GB/s

Sur les cartes haut de gammes par sur ma vieille 9600GT millieu de gamme hein.

Pour la GTX 295 ( http://www.fudzilla.com/content/view/12774/40/ ) :

GPUZ confirms the 720MHz clocks and the 242GB/s bandwidth. Note that the reference card runs at 576MHz and has a
223GB/s bandwidth.

Il faudrait en récupérer une pour tester la bande passante host <-> device.

-Donc pinned mémory, quelles différences/limitations/intérêts (à part la manière de réserver la RAM) entre les deux modes ?

Je pense que ce paper t'aidera :
http://forums.nvidia.com/index.php?a...=post&id=12274

- Pour OpenCL, utilise t'il un logiciel assembleur dédiée ?

Je crois bien que oui, mais je ne suis pas sur a 100%, si quelqu'un utilise OpenCL merci de confirmer.

je crois je jeux "Crysis" à bénéficié de cette technologie

Je ne pense pas, coder un morceau du jeu avec CUDA aurait coute chere a Crytek et on verrait de grandes differences de perf entre des carte nvidia et ati de même capacité...

Mouvement des nuages avec réflexions des rayons lumineux en provenance du soleil

Les deux exemples que tu proposes ne sont pas adapte a un environnement CUDA... Tu cherche du rendu, CUDA n'est pas fait pour ca (GPGPU = GPUs for Genetic and Evolutionary Computation). Pour ce que tu proposes il est bien plus intéressant d'utiliser des shaders...

[Rémi Coquet]

@dourouc05:

Quick Mode
Device to Device Bandwidth

lis mieux... là il est question de faire transiter le flux d'un device à l'autre.

@yamashi:

Sur les cartes haut de gammes par sur ma vieille 9600GT millieu de gamme hein.

Bravo NVIDIA, je me réjouis des horizons que cela nous ouvre pour la gestion réseau.

[HS](rien à voir: j'espère que pour ton serveur tu travailles en multicast ?)[/HS]

Il faudrait en récupérer une pour tester la bande passante host <-> device.

Pour les raisons expliquées ailleurs... ne compte pas sur moi pour te l'offrir !!!

Je pense que ce paper t'aidera

Super, la nuit s'annonce longue...

Je crois bien que oui, mais je ne suis pas sur a 100%, si quelqu'un utilise OpenCL merci de confirmer

En fait, ils sont un peu obligés, s'ils veulent être compatibles avec tout le monde... les electroniques étants différentes, il faut bien générer le code idoine.


Restent :

- La compression hard (H264) n'est-elle accessible que via CUDA ?

- Permet t'elle de travailler avec une key frame par frame etc. ou est-ce que les différents mode sont "gelés"

[yamashi]

- La compression hard (H264) n'est-elle accessible que via CUDA ?

- Permet t'elle de travailler avec une key frame par frame etc. ou est-ce que les différents mode sont "gelés"

Mais si on te réponds pas je suppose que c'est parce que personne ne sait ^^ .
Personnellement je ne travaille sur ce genre de chose du tout donc je ne suis pas en mesure de t'aider...

j'espère que pour ton serveur tu travailles en multicast ?

Pas vraiment, je fais du p2p sur un VPN intégré au serveur et aux clients ^^

Pour les raisons expliquées ailleurs... ne compte pas sur moi pour te l'offrir !!!

Hahaha, je vais bien réussir a en trouver une au boulot :p ou je la ferais acheté pour le "travail".

Sinon je pense que je vais faire un petit tuto CUDA donc si y a des intéressés, dites moi ce que vous aimeriez voir, si vous êtes intéressé par une approche théorique ou pratique...

[Rémi Coquet]

Mais si on te réponds pas je suppose que c'est parce que personne ne sait...

Le pire, c'est que ceux qui ne savent pas répondent la plupart du temps +
Je préfère nettement l'abstinence aux réponses à l'américaine...

Pas vraiment, je fais du p2p sur un VPN intégré au serveur et aux clients

Quelque chose de particulier t'a guidé vers ce choix là ? (nous travaillons sur quelque chose de semblable, où les gus se branchent et se débranchent à l'arrache et ou les flux doivent-être réduits au maximum (vidéo, audio, voip, automation etc...)).

Hahaha, je vais bien réussir a en trouver une au boulot :p ou je la ferais acheté pour le "travail"

Tiens, ça me rappel un échange récent avec Alcatiz...
A ton super travail, il en tombe beaucoup du camion ? Ce dernier pourrai verser près de chez nous, se serait très sympa, et Robin des bois pourrait nous faire un don

Sinon je pense que je vais faire un petit tuto CUDA donc si y a des intéressés, dites moi ce que vous aimeriez voir, si vous êtes intéressé par une approche théorique ou pratique...

Personnellement, quelques images pour la théorie et la circuiterie, quelques explications sur l'assembleur utilisé, et grosse part à la pratique qui permet de s'approprié par l'expérience (bout de code exemple) la technologie ciblée.

[yamashi]

Voici un benchmark de mon système permettant de traiter les clients d'un jeu sur le GPU contre le même algorithme sur le CPU.
Vous allez voir a quel point c'est intéressant.
Je ferais un paper plus tard pour expliquer comment je fais et je rendrais les sources open source pour que tout le monde puisse voir de quoi il s'agit.
A oui j'utilise un algorithme de recherche linéaire, en fait dans cet exemple on cherche a savoir si le joueur X voit les autres joueurs donc je test sous forme de zone 2D si le joueur voit ou pas les autres, donc l'algorithme est d'après mon peu de connaissance en algorithmique a évolution carre ?
Voila l'algo :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
for(int i = 0; i < max; ++i)
	{
		CSession&	working	= sessionList.at(i);;
		Zone&		z		= working.getZone();
 
		for(it = sessionList.begin();it != sessionList.end();++it)
		{
			CSession& tmp = *it;
			if(&tmp != &working)
			{
				{
					Zone& tmpZ = tmp.getZone();
					if((tmpZ.x + 2) > z.x && (tmpZ.x - 2) < z.x)
						if((tmpZ.y + 2) > z.y && (tmpZ.y - 2) < z.y)
							working.push(&tmp);
 
				}
			}
		}	
	}

Il y a surement beaucoup d'otpimisation possible a faire, voir changer d'algorithme mais c'est juste a titre de test pour comparer avec le même algorithme exécuté sur le GPU.

Voila le résultât :



Pour 8000 clients, il faut 53ms a la carte graphique et 527ms au CPU ce qui représente une amélioration des performances de 1000% !!!
Si on ajoute GPU tick et GPU memory on atteint 79ms sur le GPU mais la latence de la carte graphique peut être optimiser en utilisant de la pinned memory dans quel cas on obtiendrais quelque chose comme 60-65ms.
Dans la GPU memory je compte le temps d'allocation de la mémoire sur le GPU, le temps de la copie du vector dans un tableau C, le temps de la copie du tableau C vers le GPU, la copie des résultats vers le CPU, le renvoie des données dans le vector a partir du tableau C et la génération des clients.
Donc il n'y a pas que de la mémoire, il y a pas mal de traitement que l'on ne rencontrerait pas dans un serveur normal.

Apres avec une carte plus récente et haut de gamme, une meilleur carte mere et de la vrai ram (pas de ma noname ) les perf doivent exploser.

[Rémi Coquet]

yamashi du matin, yamashi bien... très bien même !

Est-il possible de récupérer les codes générés: assembleur CPU/GPU histoire de de comprendre tout ça mieux ?

[yamashi]

Enfin matin pour toi, soir pour moi :p (oui oui je vis aux US)
Je pense qu'il est possible de recuperer le code PTX mais je ne saurais te dire comment faire ^^

[Rémi Coquet]

Pour ce qui est des améliorations de ton algo, c'est sûr que coder ça directement en assembleur (décomptes avec "test" et "cmove"s pour éviter les saut chronovres) serait nettement plus efficace. Je n'ai pas regardé de près l'optimisation stratégique mais il faudrait remettre ça dans le contexte. SSEx devrait permettre de désynchroniser le tout pour que tes autres tâches ne soient absolument pas perturbées (quoique, tu précisais plus haut que le CPU ne gérait que le serveur .| donc...)

[yamashi]

c'est sûr que coder ça directement en assembleur

Ne pense pas être plus malin que le compilateur, les compilateurs sont de nos jours impossible a battre ou alors il faut être un monstre d'assembleur ce que je ne suis pas et je pense pas qu'on trouvera quelqu'un capable de battre un compilateur ici.
Il est vrai de dire ici qu'avec les options de compilation que j'utilise que le code compile est le plus rapide possible après quand je parle d'optimisation de l'algorithme, je parle de l'algorithme en lui même, pas du code, il faudrait revoir mon algorithme de recherche.

Note : Je n'ai pas optimise l'algorithme du CPU, mais je n'ai pas optimise celui du GPU non plus, la boucle de test et exactement la même.

[Rémi Coquet]

Ne pense pas être plus malin que le compilateur, les compilateurs sont de nos jours impossible a battre ou alors il faut être un monstre d'assembleur ce que je ne suis pas et je pense pas qu'on trouvera quelqu'un capable de battre un compilateur ici.
Il est vrai de dire ici qu'avec les options de compilation que j'utilise que le code compile est le plus rapide possible après quand je parle d'optimisation de l'algorithme, je parle de l'algorithme en lui-même, pas du code, il faudrait revoir mon algorithme de recherche.

Le malin est un menteur, le père même du mensonge et le principe du mensonge c'est la virtualisation de choses qui n'existent pas , certes...

le code compile est le plus rapide possible

Ça par contre, c'est un "vrai" mensonge temporaire comme dirait Abrash car il "est le plus rapide possible" tant que personne ne la regardé de plus près.

Je ne partage donc absolument pas tes suppositions péremptoires et m'y oppose même avec une fermeté absolue et inflexible (ho punaise !):
Je le démontre tous les jours dans mon travail, je ne suis pas le seul, et cela n'a rien à voir avec une monstruosité quelconque comme tu sembles vouloir l'idolâtrer mais avec la logique même de notre système cognitif:
Un compilateur génère un code absolument "dégueulasse" (même celui d'Intel qui est quasiment le plus propre, il faut arrêter l'intox avec ça "basta!").
C'est le discours de tous ceux qui ne font pas vraiment d'assembleur, justement, et qui transfert leur pouvoir et (espèrent) la charge de travail à un automate stupide écrit par des hommes comme les autres.

Fais-moi passer ton bout d'algo compilé avec tes supers options... je te le retourne simplement dé-assemblé et commenté avec la version ré-écrite par un enfant... histoire de clarifier les choses et de montrer que l'esprit humain de base (avec une main dans le dos) est l'auteur de toutes ces technologies et non pas leur valet.

Là où nous sommes par contre en plein accord, c'est sur le fait que l'optimisation stratégique précède à et prévaut sur l'optimisation millimétrée: On ne le répétera jamais assez.
Ceci dit, dans la stratégie, entre le fait de tout coder en assembleur afin que la stratégie soit adaptée au processeur et non adaptée au compilateur (CQFD) si si.

Si tu relis le post précédant, tu verras que tous les sauts sont supprimés et les comparaisons remplacées par des "test" (donc non-attribué), le décompte supprime la double comparaison à chaque itération et pour chaque boucle… dans un algo à haut niveau de répétition comme le tient, pas besoin de mesurer, c'est déjà plus rapide sur le papier et le compilateur ne peut pas voir ces possibilités (impossible comme tu dis).

[yamashi]

Enfin bon, assembleur humain ou machine ça ne changera pas les 1000% d'accelleration obtenu ou du moins pas de beaucoup, surtout que si on le retape sur le CPU, on le retape sur le GPU.
Sinon je pense a un autre algo, je vais le tester je vous tiens au courant.

[Rémi Coquet]

Comparons ce qui peut l'être, cad le code avec du code, pas du code sur deux électroniques différentes !
Je comprends le gain et je me réjouis vraiment avec toi de cette belle performance mais il serait très dommage de ne pas gagner ce qui est à gagner au niveau du travail CPU par un code rationel et plus efficace.
la partie client/serveur peut faire de très grands bons (je parle par expérience sur notre produit) dans des proportions qui sans atteindre totalement celles que tu annoncent n'ont pas à rougir pour le moins.