Crunchez vos adresses URL
|
Rejoignez notre discord
|
Hébergez vos photos
Page 20 sur 22 PremièrePremière ... 101213141516171819202122 DernièreDernière
Affichage des résultats 571 à 600 sur 658

Discussion: Archi GPU et GPGPU

  1. #571
    Goto a quelques slides sorti de la présentation sur CUDA 5 aussi :
    http://pc.watch.impress.co.jp/docs/c...17_533500.html

    Edit: et le whitepaper officiel : http://www.nvidia.com/content/PDF/ke...Whitepaper.pdf

    Peu à peu, les principales limitations de CUDA sautent. Avec le parallélisme dynamique et la compilation séparée, CUDA devient enfin composable : on va pouvoir écrire des bibliothèques génériques réutilisables, qui font appel à d'autres bibliothèques. Pour moi c'est indispensable à la formation d'un « écosystème » en CUDA.

    Le parallélisme dynamique signifie aussi qu'on va déporter de la logique de contrôle côté GPU. Qui sait, ça pourrait motiver l'inclusion de cores SM optimisés latence avec le même jeu d'instruction que les SMX dans les futurs GPU... Pour faire des architectures hétérogènes avec différents types de SM (SMX, SMS, SML, SMXXL...)
    Dernière modification par Møgluglu ; 17/05/2012 à 18h09.

  2. #572
    Citation Envoyé par Møgluglu Voir le message
    Goto a quelques slides sorti de la présentation sur CUDA 5 aussi :
    http://pc.watch.impress.co.jp/docs/c...17_533500.html

    Edit: et le whitepaper officiel : http://www.nvidia.com/content/PDF/ke...Whitepaper.pdf

    Peu à peu, les principales limitations de CUDA sautent. Avec le parallélisme dynamique et la compilation séparée, CUDA devient enfin composable : on va pouvoir écrire des bibliothèques génériques réutilisables, qui font appel à d'autres bibliothèques. Pour moi c'est indispensable à la formation d'un « écosystème » en CUDA.

    Le parallélisme dynamique signifie aussi qu'on va déporter de la logique de contrôle côté GPU. Qui sait, ça pourrait motiver l'inclusion de cores SM optimisés latence avec le même jeu d'instruction que les SMX dans les futurs GPU... Pour faire des architectures hétérogènes avec différents types de SM (SMX, SMS, SML, SMXXL...)
    Pour le jeu d'instructions je ne sais pas, mais avoir quelques SM optimisés latence et beaucoup optimisés débit est le but avoué de Bill Dally. Minimiser les mouvements de données, tout ça tout ça.

    T'es pas à la GTC pour voir tout ça en direct ?
    Mon blog (absolument pas à jour) : Teχlog

  3. #573
    Citation Envoyé par Alexko Voir le message
    Pour le jeu d'instructions je ne sais pas, mais avoir quelques SM optimisés latence et beaucoup optimisés débit est le but avoué de Bill Dally. Minimiser les mouvements de données, tout ça tout ça.
    Oui, mais le jeu d'instruction ça change tout. Parce que si c'est juste pour intégrer des cœurs CPU et des cœurs GPU dans la même puce, pas besoin d'attendre 2017.
    Indépendamment de NVidia, les archis hétérogènes avec cœurs OoO et cœurs SIMT à jeu d'instruction unique font totalement partie de mon plan de domination du monde à moi.

    T'es pas à la GTC pour voir tout ça en direct ?
    Déjà que je vais devoir payer de ma poche pour aller à ISCA, je vais pas en plus me déplacer juste pour assister au show de JHH et à des sessions publicitaires (et dont les vidéos seront disponibles sur l'interweb).

  4. #574
    Citation Envoyé par Møgluglu Voir le message
    Oui, mais le jeu d'instruction ça change tout. Parce que si c'est juste pour intégrer des cœurs CPU et des cœurs GPU dans la même puce, pas besoin d'attendre 2017.
    Certes, mais si c'est le même jeu d'instruction, le choix du type de SM se fait forcément en hardware, non ? Avec Kepler, NVIDIA semble vouloir prendre la direction opposée. Et quoi qu'il en soit, l'ISA des SM optimisés latence (comme Dally les appelle) sera forcément étudiée pour être aussi complémentaire que possible de l'ISA des SM débit. On ne peut pas franchement en dire autant du x86-64 des cores CPU de Trinity, par exemple. Donc c'est pas tout à fait pareil.

    Citation Envoyé par Møgluglu Voir le message
    Déjà que je vais devoir payer de ma poche pour aller à ISCA, je vais pas en plus me déplacer juste pour assister au show de JHH et à des sessions publicitaires (et dont les vidéos seront disponibles sur l'interweb).
    Forcément, vu comme ça… :D
    Mon blog (absolument pas à jour) : Teχlog

  5. #575
    Citation Envoyé par Alexko Voir le message
    Certes, mais si c'est le même jeu d'instruction, le choix du type de SM se fait forcément en hardware, non ? Avec Kepler, NVIDIA semble vouloir prendre la direction opposée.
    Je n'ai pas cette impression. Avec Echelon ou ce que semble être Denver, oui, mais dans Kepler le code de contrôle tourne sur les mêmes SMX que le code de calcul. Et le scheduling sur les SMX disponibles est fait entièrement en hardware. Donc je ne vois rien qui empêcherait d'avoir plusieurs types de SM et des décisions de placement dynamiques assistées en hardware. Bien sûr, il est peu probable que ça soit la stratégie de Nvidia. (Ce qui fait une bonne raison pour s'y intéresser. )

    Et quoi qu'il en soit, l'ISA des SM optimisés latence (comme Dally les appelle) sera forcément étudiée pour être aussi complémentaire que possible de l'ISA des SM débit. On ne peut pas franchement en dire autant du x86-64 des cores CPU de Trinity, par exemple. Donc c'est pas tout à fait pareil.
    Je ne sais pas exactement quel influence l'ISA peut avoir sur l'optimisation latence/débit d'un cœur, mais je soupçonne qu'elle n'est pas si grande. Dans l'absolu, un processeur SIMT x86 ne me parait pas plus choquant qu'un superscalaire out-of-order x86.

  6. #576
    Citation Envoyé par Møgluglu Voir le message
    Je n'ai pas cette impression. Avec Echelon ou ce que semble être Denver, oui, mais dans Kepler le code de contrôle tourne sur les mêmes SMX que le code de calcul. Et le scheduling sur les SMX disponibles est fait entièrement en hardware. Donc je ne vois rien qui empêcherait d'avoir plusieurs types de SM et des décisions de placement dynamiques assistées en hardware. Bien sûr, il est peu probable que ça soit la stratégie de Nvidia. (Ce qui fait une bonne raison pour s'y intéresser. )
    Ce que je voulais dire, c'est que si on a le même ISA pour les SM-latence (SML) et SM-throughput (SMT) alors quand le GPU doit exécuter le code, il est bien obligé de déterminer lui-même où l'exécuter. C'est possible, mais est-ce possible avec une consommation et un taux de réussite raisonnables ? Je n'en ai honnêtement aucune idée, mais ça me semble être un problème difficile.

    Citation Envoyé par Møgluglu Voir le message
    Je ne sais pas exactement quel influence l'ISA peut avoir sur l'optimisation latence/débit d'un cœur, mais je soupçonne qu'elle n'est pas si grande. Dans l'absolu, un processeur SIMT x86 ne me parait pas plus choquant qu'un superscalaire out-of-order x86.
    Oui mais si les deux ISA sont développés en parallèle et spécifiquement pour des SML et SMT, on peut avoir un bijection sémantique parfaite entre les deux, ce qui devrait simplifier un peu les choses. De cette façon, on pourrait faire le choix du SML ou SMT à la compilation, soit en se basant sur des directives du programmeur, soit sur des heuristiques automatiques qui resteraient transparentes pour le programmeur (voire les deux). Mais ça a l'inconvénient d'être statique, et quelque chose me dit que ça ne te plaît pas des masses. ^^

    Cela dit, même ISA ou pas, les limitations au niveau du partage de la mémoire sur Trinity sont telles qu'on est loin de pouvoir faire ça.
    Mon blog (absolument pas à jour) : Teχlog

  7. #577
    Citation Envoyé par Alexko Voir le message
    Ce que je voulais dire, c'est que si on a le même ISA pour les SM-latence (SML) et SM-throughput (SMT) alors quand le GPU doit exécuter le code, il est bien obligé de déterminer lui-même où l'exécuter. C'est possible, mais est-ce possible avec une consommation et un taux de réussite raisonnables ? Je n'en ai honnêtement aucune idée, mais ça me semble être un problème difficile.
    Oui, c'est pour ça que c'est intéressant.
    En fait je vois ça dans l'autre sens : grâce au fait de partager la même ISA, on peut décider dynamiquement où exécuter le code. Et si on se permet en plus de migrer des tâches au vol, façon ARM big.LITTLE, ça devient encore plus drôle.

    Oui mais si les deux ISA sont développés en parallèle et spécifiquement pour des SML et SMT, on peut avoir un bijection sémantique parfaite entre les deux, ce qui devrait simplifier un peu les choses. De cette façon, on pourrait faire le choix du SML ou SMT à la compilation, soit en se basant sur des directives du programmeur, soit sur des heuristiques automatiques qui resteraient transparentes pour le programmeur (voire les deux). Mais ça a l'inconvénient d'être statique, et quelque chose me dit que ça ne te plaît pas des masses. ^^
    Ça peut marcher, mais est-ce que ça change vraiment quelque chose par rapport à la situation d'avant le Dynamic Parallelism, où on avait du code CPU et du code GPU séparés ?…

    Mais si, je tolère ce qui est statique… tant que ça sert juste à faire remonter des infos pour guider une politique dynamique.

  8. #578
    Citation Envoyé par Møgluglu Voir le message
    Oui, c'est pour ça que c'est intéressant.
    En fait je vois ça dans l'autre sens : grâce au fait de partager la même ISA, on peut décider dynamiquement où exécuter le code. Et si on se permet en plus de migrer des tâches au vol, façon ARM big.LITTLE, ça devient encore plus drôle.
    Quel pourrait être l'intérêt de migrer des tâches au vol ?


    Citation Envoyé par Møgluglu Voir le message
    Ça peut marcher, mais est-ce que ça change vraiment quelque chose par rapport à la situation d'avant le Dynamic Parallelism, où on avait du code CPU et du code GPU séparés ?…

    Mais si, je tolère ce qui est statique… tant que ça sert juste à faire remonter des infos pour guider une politique dynamique.
    Non, fondamentalement ça ne change rien, juste la granularité de la séparation.
    Mon blog (absolument pas à jour) : Teχlog

  9. #579
    Citation Envoyé par Alexko Voir le message
    Quel pourrait être l'intérêt de migrer des tâches au vol ?
    Par exemple, tu as du code de ce genre (tiré du CUDA prog guide)
    Code:
    #include <stdlib.h>
    __global__ void mallocTest()
    {
        __shared__ int* data;
        // The first thread in the block does the allocation
        // and then shares the pointer with all other threads
        // through shared memory, so that access can easily be
        // coalesced. 64 bytes per thread are allocated.
        if (threadIdx.x == 0)
           data = (int*)malloc(blockDim.x * 64);
        __syncthreads();
        // Check for failure
        if (data == NULL)
            return;
        // Threads index into the memory, ensuring coalescence
        int* ptr = data;
        for (int i = 0; i < 64; ++i)
            ptr[i * blockDim.x + threadIdx.x] = threadIdx.x;
        // Ensure all threads complete before freeing
        __syncthreads();
        // Only one thread may free the memory!
        if (threadIdx.x == 0)
            free(data);
    }
    int main()
    {
        cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
        mallocTest<<<10, 128>>>();
        cudaDeviceSynchronize();
        return 0;
    }
    Le malloc et le free sont du code scalaire sur le chemin critique. Bon, dans ce cas, ça ne vaut probablement pas le coup de migrer. C'est plutôt un datapath scalaire optimisé latence dans chaque SM qu'il faudrait. (On y revient toujours...)

  10. #580
    Si tu partages un espace memoire coherent entre ton core optimise latence et ton core optimise throughput, tu n'as pas vraiment besoin de "migration" pour faire des malloc et des free efficacement.

    Vu que tu payes ton billet pour venir, j'offrirai la biere alors .
    fefe - Dillon Y'Bon

  11. #581
    Citation Envoyé par fefe Voir le message
    Si tu partages un espace memoire coherent entre ton core optimise latence et ton core optimise throughput, tu n'as pas vraiment besoin de "migration" pour faire des malloc et des free efficacement.
    En fait, je cherchais un exemple de portion de code scalaire dans un kernel GPGPU. Mais le malloc n'est pas tellement convaincant...

    Vu que tu payes ton billet pour venir, j'offrirai la biere alors .
    Merci.

  12. #582
    Pfff ce forum est tout mort.

    DGEMM :
    Xeon Phi SE10P : 829 GF/s pour 195W (peak 1073 GF/s)
    K20X : 1220 TF/s pour 235W (peak 1330 GF/s)

    Donc sur DGEMM meilleure efficacite pour nVidia a la fois en terme de GF/W et en terme de resultat / peak.

  13. #583
    Intéressant. Ils n'ont pas l'air d'en avoir vendu beaucoup d'ailleurs : 7 systèmes dans le nouveau Top500 contre 50 pour Nvidia. La plus grosse installation Xeon Phi (Stampede) fait à peine mieux que celle basée sur Fermi il y a 2 ans (Tianhe-1A).

    Mais c'est du x86, tavu.

  14. #584
    Citation Envoyé par newbie06 Voir le message
    Pfff ce forum est tout mort.

    DGEMM :
    Xeon Phi SE10P : 829 GF/s pour 195W (peak 1073 GF/s)
    K20X : 1220 TF/s pour 235W (peak 1330 GF/s)

    Donc sur DGEMM meilleure efficacite pour nVidia a la fois en terme de GF/W et en terme de resultat / peak.
    Je ne suis pas entièrement convaincu que K20X soit un vrai produit, par opposition à une vitrine technologique pour Titan. Notamment, le site de Cray et la brochure en PDF du XK7 ne mentionnent que le K20, moins performant et moins efficace en FLOPS/W.
    http://www.cray.com/Products/XK/Specifications.aspx
    http://www.cray.com/Assets/PDF/produ...K7Brochure.pdf
    Mon blog (absolument pas à jour) : Teχlog

  15. #585
    Citation Envoyé par Alexko Voir le message
    Je ne suis pas entièrement convaincu que K20X soit un vrai produit, par opposition à une vitrine technologique pour Titan. Notamment, le site de Cray et la brochure en PDF du XK7 ne mentionnent que le K20, moins performant et moins efficace en FLOPS/W.
    http://www.cray.com/Products/XK/Specifications.aspx
    http://www.cray.com/Assets/PDF/produ...K7Brochure.pdf
    Oui il faudrait avoir les chiffres pour le K20, mais note bien que les 235W sont le TDP de la carte alors que les 195W sont une mesure. Donc je pense que le K20 restera plus efficace de toute maniere.

    ---------- Post added at 16h49 ---------- Previous post was at 16h47 ----------

    Citation Envoyé par Møgluglu Voir le message
    Intéressant. Ils n'ont pas l'air d'en avoir vendu beaucoup d'ailleurs : 7 systèmes dans le nouveau Top500 contre 50 pour Nvidia. La plus grosse installation Xeon Phi (Stampede) fait à peine mieux que celle basée sur Fermi il y a 2 ans (Tianhe-1A).
    Xeon Phi n'est meme pas vraiment commercialise non plus, faut lui laisser le temps.

    Mais c'est du x86, tavu.
    Oui et il n'est meme pas ADV

  16. #586
    Citation Envoyé par newbie06 Voir le message
    Oui il faudrait avoir les chiffres pour le K20, mais note bien que les 235W sont le TDP de la carte alors que les 195W sont une mesure. Donc je pense que le K20 restera plus efficace de toute maniere.
    Ah oui, effectivement. Reste à voir ce que ça donne sur quelque chose de plus compliqué qu'une multiplication de matrices.
    Mon blog (absolument pas à jour) : Teχlog

  17. #587
    Citation Envoyé par Alexko Voir le message
    Ah oui, effectivement. Reste à voir ce que ça donne sur quelque chose de plus compliqué qu'une multiplication de matrices.
    Oui DGEMM reste beaucoup trop simpliste et a du etre optimise a mort par nVidia. D'un autre cote, je pense qu'Intel a fait de meme pour Xeon Phi et du coup je suis etonne qu'il soit si loin du max theorique du CPU.

  18. #588
    Citation Envoyé par Alexko Voir le message
    Ah oui, effectivement. Reste à voir ce que ça donne sur quelque chose de plus compliqué qu'une multiplication de matrices.
    Et de ce qu'Intel va proposer comme compilateur avec auto-vectorisation. Parce que s'il faut réécrire tout le code en intrinsics avec des petits bouts de vecteurs partout, ça risque de pas le faire face à CUDA.

    En résumé, on a :
    Code:
              K20       K20X        Phi 3100    Phi 5110P
    Tflops DP 1.17      1.31        1.0         1.01
    GB/s      208 (182) 250 (219)   240 (233)   320 (310)
    GB Mem    5 (4.4)   6 (5.25)    6 (5.8)     8 (7.75)
    TDP       225       235         300         225
    Le Xeon Phi est un monstre de bande passante. Les chiffres entre parenthèse sont avec l'ECC activé.

  19. #589
    Citation Envoyé par Møgluglu Voir le message
    Et de ce qu'Intel va proposer comme compilateur avec auto-vectorisation. Parce que s'il faut réécrire tout le code en intrinsics avec des petits bouts de vecteurs partout, ça risque de pas le faire face à CUDA.
    Attends ils ont icc et avec une ligne ca fait tout tout seul, la preuve : meme des journalistes ont pu le faire. Je ris encore de la naivete de Charlie

  20. #590
    Mon dieu, on dirait les pubs pour OpenACC.

    Et ils prétendent vraiment faire tourner le binaire compilé pour Xeon Phi sur un Xeon, ou c'est Charlie qui a loupé un truc ?

  21. #591
    Citation Envoyé par Møgluglu Voir le message
    Et ils prétendent vraiment faire tourner le binaire compilé pour Xeon Phi sur un Xeon, ou c'est Charlie qui a loupé un truc ?
    Charlie (par l'intermediaire de Copper) pretend qu'Intel a confirme que ce serait bien le meme binaire, ce qui est impossible sauf si le code utilise du x87, et la, comment dire, je doute

    EDIT : Remarque ils ont ptet mis du code manage...

  22. #592
    320GB/s... Oh la vache...
    Par contre, si y'a pas de sortie vidéo, c'est un peu dommage pour ce que je voulais faire (calcul ultra-rapide sur GPU et affichage sur place).

  23. #593
    77% peak sur DGEMM est effectivement un peu bas pour le Phi, typiquement les x86 des 10 dernieres annees atteignent 85% a 90% du peak.
    L'efficacite du G20X est par contre particulierement impressionante.
    fefe - Dillon Y'Bon

  24. #594
    D'un autre côté, le GK110 est probablement "fait pour". Les 256 registres + l'instruction swizzle doivent bien aider dans DGEMM. Jusqu'au GK104, les facteurs limitants étaient la taille de blocking limitée par le nombre de registres, et le débit de la shared memory.
    Ça reste bien sûr un accomplissement d'arriver à tout faire bien marcher pour que ça passe à l'échelle…

  25. #595
    Question que je me pose pour un dossier :

    Je compare Tegra 3 aux autres GPU mobiles. Tegra 3 a un rendu type PC (immédiat), les autres sont en tile avec du HSR (PowerVR) ou du « early z ».

    Dans les jeux optimisés Tegra 3, les développeurs font usage essentiellement de deux choses : des effets de glow sur les objets et des lumières plus travaillées. Qu'est-ce qui fait que ce genre de choses est plus rapide/possible sur Tegra 3 et pas sur les autres GPU ?

    De ce que j'ai pu comprendre, ça vient surtout du fait que Tegra 3 a 8 unités de pixels et 4 unités de Vertex.

    Autre question, j'ai vu plusieurs fois que les GPU en tile avec HSR souffrait plus d'une augmentation de la géométrie de la scène, mais pourquoi ? Parce qu'augmenter la géométrie demande plus de boulot au niveau de l'analyse pour le HSR ?

  26. #596
    Faudrait creer un thread special pour Intel et les GPU.

    Apres ses differents designs, apres ImgTec, apres LRB, ZiiLabs, ex 3dlabs.

    http://www.techweekeurope.co.uk/news...to-intel-99619

    C'est du delire l'histoire d'Intel et du GPU

  27. #597
    La GTX Titan recupere enfin de la performance sur les DP. Dommage ca coute 1000 Euros

  28. #598
    C'est déjà nettement moins qu'un Tesla K20... Moins d'un euro le gigaflop, c'est raisonnable.

  29. #599
    NVIDIA rachète PGI à ST.
    http://blogs.nvidia.com/blog/2013/07/29/portland/

    Ça a du sens. Bientôt un front-end OpenACC dans CUDA ?

  30. #600
    Ou une maniere d'offrir d'autres debouches aux ingenieurs de la region de Portland ?
    fefe - Dillon Y'Bon

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
  •