Hello,
 
Je me pose certaines petites questions à force de lire des tutos sur le net, j'ai la vague impression que parfois certains se contredisent...
Si vous trouvez le temps de répondre à une ou deux, je vous en serais très reconnaissant... ! (je pourrai éditer ce post pour faire un résumé pour que ça puisse servir à d'autres!)
 
- A-t-on bien, à chaque instant, une seule grille (grid) d'active sur une carte ? Peut-on avoir plusieurs grilles s'éxécutant en parallèle (dans un programme multi thread-CPU par exemple) ?
lorsque dans mon code host j'appelle 2 kernels à la suite, l'appel au second attent-il que le premier ait fini, ou faut il que je synchronise le tout à la main si besoin est ?
- A-t-on en permanence au maximum un seul warp actif (en cours d'exécution) par bloc ?
- J'ai pu lire par endroits un découpage en demi-warp : les 2 demis-warps s'exécutent-ils en parallèle ou en série ? Je parierai sur le "série", puisque apparamment il ne peut pas y avoir de conflits de banques entre deux demis warps
- Pourquoi y a t il un découpage en demis warps ? Est-ce une contrainte hardware/technique, ou y-a-t-il une raison plus profonde ?
- Comment déterminer quand l'exécution que je demande va être sérialisée à cause du découpage de la grille que j'ai fait ?
par exemple, si j'appelle des kernels avec <<< 50, 32 >>>, est-ce que le nombre 50, qui est, en l'occurence, supérieur au nombre de multiprocesseurs sur ma carte (30), est-ce un mal ?
De même, si je fais <<< 16, 64 >>>, y-a-t-il sérialisation de l'exécution des warps ? (2x32 = 64) (en fait, ça revient à ma question 2)
- Au cas ou j'ai raconté des bêtises, quel rôle les multiprocesseurs jouent-ils dans un programme CUDA ? Centralisent-ils l'exécution d'une entité particulière, ou est-ce que ça n'a rien à voir ?
Plus précisément, comment prendre en compte intelligemment dans le code le nombre de "threads processors" et de "multiprocessors" ?
 
Un dernier truc : on est bien d'accord que sur http://jeux.developpez.com/faq/gpg [...] ocabulaire
Les définitions de grille et blocs sont bien "à l'envers" ?
 
Si vous avez de bonnes références (suffisamment exhaustive/technique) autres que la doc NVIDIA (qui est quand même pas mal du tout je trouve) , je suis preneur. L'article "CUDA approfondi" de developpez.com m'a l'air pas mal, de mon point de vue de débutant, mais quand je lis le dernier post de http://www.developpez.net/forums/d [...] sons-cuda/ j'ai l'impression de passer à côté de certaines choses
 
Merci pour le moindre coup de pouce !

la doc nvidia est a chier niveau hardware. Je te conseille les papiers de Volkov sur la programmation des GPU en mode SIMD qui est beaucoup plus efficace.

Si tu appelles 2 kernels depuis 1 thread host, ils seront sequentialisé, tu peux lancer 2 threads host, chacun lancant un kernel sur une sous grille.
 

?? Je comprends pas. Les warps actifs sont ceux necessaires à couvrir tes besoisn en block/threads
 

le demi-warp es l'unité matérielle la plus pêtite ne necessitant pas de synchro interne. Donc es demi warp s'executent en parallele
 

contrainte hard
 

En gros, si tu demande plus de bloc et de thread que t'as carte  de multi-core physique, il y a serialization.
Cette serialization est trés couteuse et en general on evite.
 

Ce n'est pas le bon modèle. NVIDIA a pondu ça à la va vite. Le bon modele est de voir la GPU comme un ensemble de machine SIMD synchronizable.
 

JV.com what else :€

 
Pour te chiffrer la réponse de Joel,
D'après les specs :  
 
The maximum number of active blocks per multiprocessor is 8;
The maximum number of active warps per multiprocessor is 32;
 
Tu as 30 multiprocs, donc 30*8=240 blocks actifs simultanément
OU
32*32=1024 threads actifs par multiproc, c'est à dire 1024*30=30720 threads actifs (dans le sens où ils l'entendent...)
 
Si tu dépasses 240 blocs, ça sérialise, si tu dépasses 30720 threads, ça sérialise
 
 

c'est pas 32*8 ?

je sais pas, le monsieur il dit qu'il a 30 multiprocs.
 
Les warps font 32 threads, donc 32*32*30 threads simultanés.
et 30*8 pour le nombre de blocs simultanés
 sauf si je délire

ah ok j'ai mal lu v_v

Hé bien, merci du coup de main... !

J'ai déjà étudié bien en détail deux des articles de Volkov, vraiment merci pour la référence... en fait je dois coder une EDP avec un schéma aux différences finies à coucher dehors, ça m'a donné pas mal d'idées.
Juste avant de me lancer dans la phase d'analyse puis codage, j'aurais encore quelques questions, si ce n'est pas trop demander :

- Il y a 16384 registres/multiproc, qui sont - souvent - utilisés pour les variables locales aux kernels si j'ai bien compris. Puisque ces registres sont partagés (un thread ne doit pas monopoliser tous les registres, ça serait assez idiot), j'imagine que le compilo nvcc doit déterminer à la compilation comment les registres vont être utilisés. Le hic, c'est que le compilo ne sait pas à l'avance combien de threads vont se partager les registres d'un multiproc... ! (puisque la configuration d'appel aux kernels est dynamique, sans parler du multithreading côté host qui peut - j'imagine - générer différentes activités dans un même multiproc) Adopte-t-il alors une stratégie "agressive" : ie. actuellement j'ai 1024 threads actifs maximum par multiprocesseur pour 16384 registres; le ratio valant 16, le compilo limite-t-il l'utilisation des registres à 16/thread ?
- J'ai pu lire (je ne sais plus, et je ne le retrouve pas) que les 16 kB/block de shared memory ne sont pas accessibles entièrement (sont-ils utilisés pour passer les arguments aux kernels ?). Ma question est alors : comment peut-on connaître précisément la mémoire disponible ? Et surtout, ou se trouve la zone à laquelle je n'ai pas le droit de toucher ? A la fin des 16 kB, ou je dis des bêtises ? Juste histoire d'avoir un contrôle fin des ressources disponibles.

- [EDIT - question idiote enlevée]

Voilou, en tous cas, encore merci pour votre aide qui m'a déjà été précieuse... !

En fait je me rends compte que ma première question contient un non-sens, mais ça ne change pas mon interrogation : comment nvcc détermine-t-il l'utilisation des registres à la compilation ?
Y a-t-il moyen de lui forcer la main quand c'est nécessaire ?

Pour ta question, personne ne sait (NVIDIA non plus je parie ).
 
Nous on se casse pas les pieds on lance exactement le nombre physiques de bloc et de threads par warp et on fait un do while sur les données en internes. Comme ça on maitrise tout et nvcc ne fait rien de moche.

Si le kernel qui a été compilé est incompatible avec le nombre de threads que tu demandes, il ne s'exécute pas.
 

cf doc nvidia
--ptxas-options=-v
-maxregcount N

 
Hmm, je ne comprends pas bien ce que tu dis à vrai dire :/
Qu'entends-tu par "on lance exactement le nombre physiques de bloc et de threads par warp" ? Vous avez une taille "statique" de block/grid dans tous vos programmes ? J'imagine que ça exclut aussi du multithreading ?
Et, personne ne s'est encore amusé à "rétro-ingénieurer" l'espèce d'assembleur que nvcc produit depuis le temps, pour répondre à la question ? Ca ne donne rien/c'est peut-être trop imbitable ? (youpi, j'ai trouvé de quoi m'occupper)
 
Et, vu que tu as l'air relativement critique vis-à-vis de nvidia, peut-être pourrais-tu me conseiller des documents du genre "the dark side of cuda" si ça existe ? J'ai l'impression qu'on y doit pouvoir en apprendre plus que dans le programming guide officiel :)

 
Et sa compatibilité, c'est à moi de l'assurer en m'aidant de la sortie de ptxas-options=-v donc?
 

 
Ok, j'ai peut-être mis la charrue avant les boeufs je n'ai pas encore lu les options du compilo, merci (et désolé) !
Pour le maxregcount je ne vois pas le rapport avec la quantité de shared memory par contre... !

oui si tu y arrives. sinon au pif ça marche aussi
 

 
Tu cherchais un moyen de limiter la quantité de registres utilisés. (pas testé)

Ok, bon c'est super tout ça... !
A priori j'aurai une dernière question (j'espère) dont je ne trouve pas de réponse : pour pouvoir optimiser mon calcul, j'imagine que j'aurai à jouer avec les warps/demis warps, et donc :
 
=> Supposons que mes blocs soient de taille (8, 8, 2) par exemple (128 threads/bloc), comment identifier les même threads d'un warp et d'un même demi warp ? quelle est la relation d' "ordre" entre les threads en 3D (et même en 2D)
 
Merci encore...

 
Les blocs ont un numéro, les threads dans le bloc aussi. De tête il existe un schéma comme tu demandes dans la doc.
 

Bizarre, je viens de passer 50 bonnes minutes à farfouiller un peu partout dans la doc/sur le net, je ne vois rien de tel... !
J'ai bien compris les threadIdx.x/y/z, blockIdx.x/y, blockDim.x/y/z, mais je ne vois pas ou il est écrit comment sont groupés les threads dans un warp/demi warp dans des blocs 2D ou 3D
 
J'ai trouvé des trucs alléchants ici : http://kratos.cimne.upc.es/trac/ch [...] _kernels.h
le type utilise une fonction GlobalIdx(), seulement je ne vois pas trop d'ou elle sort (elle n'est pas dans la spec cuda et il en la définit nulle part)
 
EDIT : forcément, je poste et je trouve juste après :
 

 
Source http://theinf2.informatik.uni-jena [...] 8-p-73.ppt
 
Bon, c'était "logique", mais j'ai l'impression que l'info est assez rare...

Sur le forum Nvidia, tu trouveras 1000 fois les réponses aux questions que tu te poses.
Je pense qu'on passe tous par les mêmes problèmes.

Sinon également, plus que la doc cuda, tu peux regarder les examples du sdk qui sont assez souvent accompagnés d'une doc  en pdf qui explique pourquoi l'exemple a été implémenté comme ça. Les exemples ne sont pas forcément optimaux mais sont des algos courants.