Dans le cadre de la rubrique «Poser une question à un expert Intel», nous avons demandé au principal spécialiste Intel Konstantin Vladimirov de répondre aux questions liées à la programmation hétérogène, à la
boîte à outils oneAPI et à des choses intéressantes connexes. Le résultat a dépassé toutes nos attentes. Konstantin n'a pas perdu de temps et a donné des réponses détaillées et justifiées, sans craindre d'être polémique. En fait, nous avons eu une petite conférence sur la programmation inter-architecturale sous toutes ses formes: nuances de déchargement, optimisations, normes, etc.
Nous transférons le microphone à l'expert. Eh bien, les commentaires sont donnés au public.
Question
Soarex16Dans quelle mesure la transition d'OpenCL à oneAPI sera-t-elle laborieuse et quels avantages en retirer?
La réponse. Passer à DPC ++ peut être délicat, mais à mon avis, cela en vaut la peine. Il y a deux étapes principales.
Tout d'abord, il s'agit d'une transition de votre langage de programmation hétérogène (OpenCL, Vulkan compute), qui, très probablement, est basé sur l'API. Ici, vous avez une longueur d'avance dans le fait que vous connaissez déjà le sujet, et la difficulté est de faire passer la réflexion du contrôle direct via l'API à des constructions de langage légèrement plus implicites.
Deuxièmement, il s'agit d'une transition de votre langue hôte. Si vous avez déchargé toute votre vie du C pur, le seuil d'entrée est égal au seuil de passage du C au C ++, qui est assez élevé.
Pourquoi essayer?
Premièrement, DPC ++ fait un excellent travail pour un programmeur. Vous oublierez très rapidement, comme un cauchemar, tous ces appels explicites à clXXXYYY, et ce que signifie le sixième argument, et si vous avez oublié le code retour. De nombreux wrappers orientés objet ne cachent pas la routine, mais généralement au prix du passage de l'API OpenCL standard à l'API wrapper pas si standard (j'ai également vu ces vélos). Dans le cas de DPC ++, vous écrivez simplement le SYCL standard avec des extensions Intel (qui pourraient bientôt devenir aussi le SYCL standard).
Deuxièmement, DPC ++ prévoit une compilation conjointe, c'est-à-dire que vous pouvez être sûr des types et vous n'aurez pas de problèmes aux frontières de l'API avec les dimensions, le remplissage, l'alignement. Vous écrivez le noyau et le code hôte dans un fichier, et c'est le même code. En utilisant USM, vous pouvez également travailler avec des structures de données complexes beaucoup plus facilement.
Troisièmement, DPC ++ est un vrai C ++, c'est-à-dire qu'il permet une programmation généralisée. Par exemple, le noyau le plus simple pour ajouter deux vecteurs:
auto kern = [A, B, C](cl::sycl::id<1> wiID) { C[wiID] = A[wiID] + B[wiID];
Même chose sur OpenCL:
_kernel void vector_add(__global int *A, __global int *B, __global int *C) { int i = get_global_id(0); C[i] = A[i] + B[i]; }
Vous voyez, j'ai été obligé de pointer vers un type OpenCL int. Si j'ai besoin d'un flottant, je devrai soit écrire un autre noyau, soit utiliser un préprocesseur, ou une génération de code externe. Mettre presque toutes les fonctionnalités de C ++ à votre disposition peut être un peu effrayant si vous n'avez pas d'expérience avec C ++. Mais c'est une chose courante quand il s'agit d'un changement technologique majeur.
Et tous les avantages ne se limitent pas à cela. Je mentionnerai autre chose dans les réponses suivantes.
J'aurais donc téléchargé le compilateur à votre place et je l'aurais essayé, car ce n'est pas difficile de le faire avec le package
OneAPI .
Question
JusterOpenVINO et oneAPI seront-ils liés d'une manière ou d'une autre?
La réponse. La distribution OpenVINO fait maintenant partie de la distribution OneAPI. L'apprentissage et l'utilisation des réseaux de neurones sont des tâches difficiles à calculer qui bénéficient grandement d'une programmation hétérogène. Je pense que tôt ou tard tous les composants OneAPI permettront d'utiliser toutes les ressources informatiques à votre disposition: accélérateurs graphiques et accélérateurs spéciaux comme Nervana et FPGA. Et tout cela sans quitter le paradigme du langage et le système de type de votre programme C ++.
Questions du courrier
J'essaie de comprendre à quoi ressemblera l'accélérateur matériel AI dans 3 ans, aidez-moi. Il existe une société intéressante Graphcore et son IPU - cet appareil n'est pas moins efficace que FPGA, mais il est beaucoup plus facile à programmer - Python avec le support de TensorFlow et d'autres frameworks. Il s'avère que si les promesses de Graphcore sont remplies, il n'y aura pas besoin de FPGA sur le marché de l'apprentissage automatique. Python est beaucoup plus pratique pour les datascientists que C ++.
Êtes-vous d'accord pour dire que FPGA n'est pas adapté au marché de l'apprentissage automatique par rapport aux solutions programmables Python? Si le marché ML est perdu, quelles autres applications FPGA répandues voyez-vous?
Dans quelles applications voyez-vous le besoin inévitable d'une programmation hétérogène, où vous ne pouvez pas vous en sortir avec des outils plus pratiques comme Python?
La réponse. J'ai jeté un coup d'œil sur le type d'UIP. Encore un morceau de fer sur lequel tout le monde se déchargera. Ces gars-là rivalisent avec le GPU et avec des accélérateurs spéciaux, et non avec le FPGA.
Dans les tâches pour lesquelles un matériel spécialisé est affûté, il battra toujours FPGA, par exemple, le rendu vidéo est meilleur sur une carte vidéo, etc. Mais dans le monde (y compris dans le monde du ML), il y a beaucoup de tâches pour lesquelles rien de spécial n'a été inventé ou publié, et ici le FPGA sera toujours indispensable. Par exemple, parce qu'il y a une question de prix et, pour être bon marché, un matériel spécialisé doit être massif.
Supposons maintenant que l'UIP spécifiée est vraiment cool. Cela n'annulera pas la programmation hétérogène, au contraire, la présence d'un tel excellent accélérateur le stimulera. Et cela donnera également une longueur d'avance géante à OneAPI et DPC ++, car tôt ou tard, quelqu'un dira "Je veux utiliser à la fois votre IPU et mon GPU à partir d'un seul programme." Assez tôt parce que la programmation hétérogène est à ce sujet. Sa signification est le déchargement d'une tâche appropriée vers un appareil approprié. Une tâche peut venir de n'importe où. Et cet appareil peut être n'importe quoi, il peut même être le même appareil sur lequel le programme s'exécute. Par exemple, si vous déchargez le noyau écrit en ISPC et utilisez au maximum les fonctionnalités vectorielles de Xeon, vous pouvez le décharger vous-même tout en étant un gain significatif. Le critère principal ici est la performance. Eh bien, il n'y aura jamais trop de productivité dans ce monde. Même avec les meilleurs accélérateurs du monde.
Quant à Python et sa commodité ... Je dois admettre tout de suite que je n'aime pas les langages typés dynamiquement: ils sont lents et au lieu d'une erreur de compilation normale, vous devez attendre deux heures avant de tomber dans le runtime en raison du mauvais type. Mais je ne vois pas à quel point il est mauvais de faire les mêmes déchargements sous Python. Soit dit en passant, OneAPI comprend déjà Intel Distribution for Python, ce qui est extrêmement pratique pour diverses critiques.
Autrement dit, dans le monde de rêve des amateurs de Python, vous écrivez un programme dessus et le déchargez sur tous les accélérateurs que vous pouvez trouver en utilisant OneAPI, et non pas sur un tas de bibliothèques spécifiques au fournisseur. Une autre chose est qu'avec cette approche, vous manquez de taper de bout en bout et revenez dans le monde extrêmement précaire de la programmation basée sur API. Peut-être que le développement de DPC ++ encouragera la communauté à utiliser plus activement des outils plus appropriés, tels que C ++.
Question du courrier
Performance contre OpenCL. Il doit y avoir des taxes sur le luxe - c'est-à-dire les frais généraux. Y a-t-il des mesures?
La réponse. Sur Internet, vous pouvez trouver de nombreuses mesures avec une variété de résultats, selon le compilateur, la tâche et la qualité de la mise en œuvre. À titre de recherche personnelle, j'ai mesuré sur des tâches simples (SGEMM, DGEMM) sur mon ordinateur portable (graphiques Skylake intégrés), et j'ai vu que jusqu'à présent, il y a un certain rabattement (en pourcentage). Mais il me semble que c'est une conséquence du fait que tout cela est jusqu'à présent en version bêta.
En théorie, le résultat devrait être une accélération, pas une décélération, c'est-à-dire qu'en principe tout ce luxe devrait avoir une valeur négative. Tout tourne autour du compilateur. Lorsque votre programme se compose d'une seule source et est traité comme un seul programme, le compilateur offre des possibilités d'optimisation fantastiques et incroyables: présentation de code commun, inversion de boucles, réorganisation des sections de code et tout ce que le compilateur ne peut tout simplement pas faire dans l'approche basée sur l'API, mais tôt ou tard, elle apprendra certainement avec un modèle à source unique.
De plus, DPC ++ aura un coût négatif en termes de temps de développement. Un exemple simple est les accesseurs SYCL, que le compilateur utilise déjà pour organiser les événements et gérer les files d'attente asynchrones.
deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto C = bufferC.template get_access<sycl_write>(cgh); .... deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto D = bufferD.template get_access<sycl_write>(cgh);
Ici, le compilateur voit que les deux packages ne lisent que A et B et écrivent des tampons indépendants C et D, par conséquent, il voit la possibilité de les envoyer en parallèle s'il y a suffisamment de tailles globales.
Bien sûr, un programme OpenCL écrit de manière pédante peut le faire aussi bien, mais le temps de développement passé avec un noyau non trivial ne sera pas comparable.
Question du courrier
Toutes les façons d'optimiser les applications OpenCL pour DPC ++ sont-elles pertinentes? Quoi de neuf à y ajouter?
La réponse. Je dirais que le compilateur peut et devrait faire la plupart de l'optimisation manuelle fine qui est effectuée par les auteurs du noyau. De la même manière, par exemple, je considère comme une pratique néfaste d'installer manuellement un assembleur en ligne dans des programmes C ++, car même s'il donne des avantages tactiques, il interfère avec les optimisations et agit comme un facteur négatif dans le développement et le transfert d'un produit. Eh bien, OpenCL est maintenant également assembleur.
Quant à la réponse plus détaillée, j'ai peur de l'abîme ici. Par exemple, il existe un document Intel bien connu "OpenCL Developer Guide for Intel Processor Graphics". Et il y a une
section sur la façon d'essayer, afin de ne pas mettre où la synchronisation excessive.
Donc, de mon point de vue, c'est une tâche non humaine en principe. Les gens sont extrêmement pauvres pour raisonner sur la synchronisation multi-thread et ont tendance à sculpter la synchronisation de manière conservatrice ou incorrecte, ou les deux à la fois - je mets des virgules comme ça (
mais nous l'avons corrigé - note éditoriale ).
D'un autre côté, en DPC ++, au lieu d'écrire du code avec des barrières explicites, comme ceci:
for (t = 0; t < numTiles; t++) { const int tiledRow = TS * t + row; const int tiledCol = TS * t + col; Asub[col][row] = A[globalRow * AY + tiledCol]; Bsub[col][row] = B[tiledRow * BY + globalCol];
Vous écrirez très probablement une itération explicite de
parallel_for_work_group , dans laquelle
group.parallel_for_work_item cgh.parallel_for_work_group<class mxm_kernel>( cl::sycl::range<2>{BIG_AX / TS, BIG_BY / TS}, cl::sycl::range<2>{TS, TS}, [=](cl::sycl::group<2> group) {
Par conséquent, vous n'avez pas du tout besoin de définir la synchronisation à la main et la section entière peut être supprimée.
Et vous pouvez donc marcher dans toutes les sections. Quelque chose survivra, quelque chose partira. Je prévois l'émergence d'un nouveau document «Optimisation pour DPC ++», mais le temps devrait passer, car toutes les techniques réellement efficaces ne sont développées que plus tard et avec du sang
Question du courrier
Il y a une limitation dans OpenCL - vous ne pouvez pas utiliser des "données distantes" dans le noyau, c'est-à-dire, par exemple, implémenter un "filtre large" qui utilise les données d'entrée d'un grand groupe de pixels plus grand que le groupe de travail OpenCL dans un calcul. Que propose DPC ++ à cet égard?
La réponse. Eh bien, c’est impossible. Bien sûr, je n'écris pas particulièrement les noyaux ... Mais il est absolument certain que vous pouvez utiliser toute la mémoire globale telle quelle, il vous suffit de vous assurer que vous travaillez avec les opérations atomiques (ou synchronisez les noyaux hiérarchiques en externe). Et vous pouvez également connecter le système SVM (enfin, ou l'USM en DPC ++).
Hélas, tout cela est extrêmement inefficace, et je n'aime pas toutes ces astuces. De plus, ils sont difficiles à optimiser par le compilateur.
Et donc, si nous parlons de solutions directes et efficaces, alors, bien sûr, il n'y a pas de magie dans DPC ++. Votre programme à la fin est toujours divisé en parties: le code hôte et le code de l'appareil, et toutes les restrictions de l'appareil affectent le code de l'appareil. La taille maximale du groupe de travail est ce véritable parallélisme dont votre matériel est capable. Tout cela n'est que des moyens de sortir, ce qui affecte considérablement les performances. C'est pourquoi DPC ++ offre une opportunité de le faire:
device.get_info <sycl :: info :: device :: max_work_group_size> () , puis de décider comment vivre avec le nombre résultant.
Il serait tentant, bien sûr, de créer un modèle en DPC ++, lorsque le programmeur fonctionne comme vous le souhaitez avec des boucles de n'importe quelle longueur, et que le compilateur regarde quoi faire ensuite, mais ce serait une erreur mortelle, car cela cacherait des constantes, voire des asymptotiques d'une complexité accrue l'informatique apparaissant de nulle part. Pour une autre raison, Alexandrescu a écrit que «l'encapsulation de la complexité devrait être considérée comme un crime», et cela s'applique également.
Parfois, la révision de l'algorithme lui-même est utile. Ici, DPC ++ facilite les choses car un code plus structuré est plus facile à refactoriser. Mais c'est une consolation.
Question du courrier
DPC ++ est basé sur SYCL. Mais que se passe-t-il si vous allez plus loin sous le capot, quelles sont les différences d'OpenCL dans la mise en œuvre du back-end, le cas échéant. Par exemple, le mécanisme de distribution entre les appareils hétérogènes est-il le même qu'OpenCL?
La réponse. Si vous vous mettez sous le capot, c'est OpenCL. Tous les avantages et les points forts de SYCL sont les avantages et les points forts du langage, c'est-à-dire le frontend. De la partie frontale vient le bon vieux SPIRV qui va au backend et là il est optimisé (souvent déjà au moment de l'exécution, c'est-à-dire, c'est JIT) déjà pour une carte vidéo spécifique de la même manière qu'OpenCL serait optimisé pour elle.
Une autre chose est que le mécanisme de distribution du travail entre les appareils hétérogènes est juste plus frontal que principal, car c'est le code hôte qui décide quoi envoyer et où. Et le code hôte est obtenu à partir de DPC ++. J'ai déjà montré un exemple un peu plus haut comment le compilateur peut, sur la base d'accesseurs, prendre une décision sur la parallélisation de paquets. Et ce n'est que la pointe de l'iceberg.
Question du courrier
Bibliothèques Oui, nous ne parlons pas de CUDA. Mais nous savons que pour les développeurs CUDA, il existe des bibliothèques très utiles qui fonctionnent avec de hautes performances sur le GPU. OneAPI contient également certaines bibliothèques, mais, par exemple, IPP - il n'y a rien d'archivage utile pour travailler avec des images dans oneAPI / OpenCL. Y aura-t-il quelque chose, et comment dans ce cas passer de CUDA à oneAPI?
La réponse. La transition de CUDA à un standard ouvert unique sera difficile, mais inévitable. Bien sûr, CUDA dispose désormais d'une infrastructure plus mature. Mais les caractéristiques de sa licence sont un inconvénient bloquant, car de plus en plus d'acteurs apparaissent sur le marché des systèmes hétérogènes, de plus en plus de cartes et d'accélérateurs intéressants de différents fabricants.
La diversité des API existantes rend difficile l'utilisation de ce monde de possibilités pour les programmeurs ayant une expérience du CPU classique. Ce qui conduit à OneAPI ou quelque chose comme ça. Ici, la magie n'est pas dans la percée d'Intel dans les graphiques, mais dans le fait qu'Intel ouvre la porte à DPC ++ à tout le monde. Nous ne possédons même pas la norme SYCL, elle appartient au groupe Khronos et toutes les extensions Intel sont des extensions de Khronos où tout le monde peut s'engager (et il y a des représentants de tous les principaux acteurs là-bas). Et cela signifie (les bibliothèques) et la communauté apparaîtront (apparaissent déjà), et beaucoup de postes vacants dans cette direction.
Et bien sûr, IPP sera réécrit pour de nouvelles réalités. Je n'ai rien à voir avec IPP, mais utiliser DPC ++ est du bon sens, et des gens sensés sont assis là.
Mais plus important encore, c'est maintenant le moment de l'histoire où vous pouvez écrire votre propre bibliothèque, qui dépassera IPP et que le monde entier utilisera ensuite. Parce que les standards ouverts gagnent toujours.
Question du courrier
Si nous comparons le lancement d'algorithmes de formation et d'inférence de réseaux neuronaux sur Nervana et FPGA - quelles sont les différences de programmation et l'efficacité qui en résulte?
La réponse. Je ne sais rien des détails de la programmation FPGA, j'écris des compilateurs. Mais j'ai une contre-question. Et comment allons-nous comparer? Sur les repères standard, c'est antisportif, Nervana a léché sous eux. Mais si vous avez quelque chose d'intéressant, le FPGA vous déliera les mains, et mettre ce quelque chose sur Nervana peut être long, cher, c'est tout.
Il s'avère que la question elle-même est, pour ainsi dire, issue de la série «qui est plus fort qu'un éléphant ou une baleine». Mais ce n'est pas une vraie question. La vraie question est: comment atteler à la fois un éléphant et une baleine dans une charrette? Eh bien, ou du moins distribuez, disons, un éléphant pour le tirer par terre et une baleine par mer.
Dans le cas de OneAPI, vous aurez le même programme en général en C ++ standard. Et vous pouvez l'écrire vous-même et l'exécuter avec déchargement dans les deux sens. Ce sera la tâche même qui vous intéresse, sur laquelle vous pouvez vous-même mesurer et optimiser les performances. Un standard unique et une interface unique pour les appareils hétérogènes constitueront une étape vers la comparaison des pommes avec les pommes dans ces domaines.
Par exemple: "quoi de mieux pour% de ma tâche% du point de vue de la facilité de programmation et d'efficacité - mettre cette partie sur FPGA, laisser celle-ci sur Nervana ou diviser cette partie en deux, et réécrire cette partie pour le GPU?"
Et toute l'histoire avec OneAPI - c'est juste pour vous de dire, "pourquoi y penser longtemps, je vais l'essayer maintenant rapidement, C'EST SIMPLE".
Pas encore, pas facile. Mais il y en aura.
Mot de l'expertMerci à tous pour vos questions. Il est possible et même probable que je me sois trompé, inexact et fait des erreurs. Cela arrive, sur Internet, constamment quelqu'un a tort.
J'espère que j'ai pu intéresser quelqu'un à la programmation hétérogène et au DPC ++. Je veux recommander à tous le site
sycl.tech , où se trouvent des tonnes de rapports, y compris d'experts de renommée mondiale (l'anglais est requis)
Bon à tous!
PS de l'éditeur. Cette fois, par décision unanime du comité de rédaction, il a été décidé d'attribuer le prix de la meilleure question ... à l'auteur des réponses. Je pense que vous conviendrez que c'est juste.