Comment optimiser l'utilisation des op GPU d'apprentissage profond TVM? Apprenez � vous pour atteindre 2-3 fois la mise � jour avec des dizaines de lignes de Python

Il y a quelques jours, l'�quipe Chen Tianqi a annonc� TVM, sur dit le microblogging, � nous avons publi� TVM aujourd'hui, et forment ensemble une �tude approfondie NNVM pour compl�ter une vari�t� de cha�ne d'outils d'optimisation du mat�riel pour supporter le t�l�phone mobile, cuda, OpenCL, m�tal, javascript et divers autres arri�re-plan. Bienvenue � la th�orie du compilateur apprentissage en profondeur, calcul haute performance, l'acc�l�ration mat�rielle des �tudiants int�ress�s � se joindre ensemble pour promouvoir les grands projets open source DMLC communaut� �.

Selon Lei Feng r�seau AI Technology Review est entendu que la plupart des syst�mes existants sont optimis�s pour une gamme �troite de GPU au niveau du serveur, et la n�cessit� de d�ployer beaucoup de travail sur, y compris les t�l�phones mobiles, les �quipements IOT et des acc�l�rateurs d�di�s. Et TVM est une pile mat�rielle d'extr�mit� de la profondeur du d�ploiement de la charge de travail d'apprentissage IR (repr�sentation interm�diaire). En d'autres termes, ce type de solution peut �tre distribu�e au mod�le d'apprentissage en profondeur sur une vari�t� de p�riph�riques mat�riels, pour atteindre fin � l'accord final.

Il est la pr�sence de trois caract�ristiques:

Pour optimiser CPU, GPU et autre mat�riel informatique sp�cialis� sur les t�ches d'apprentissage profondeur r�guli�re;
calcul figure peut �tre convertie automatiquement, de sorte que l'utilisation de la m�moire est r�duite au minimum, afin d'optimiser la pr�sentation des donn�es, le mode de calcul de la fusion.
Compil� � partir de la fin de l'avant existante pour fournir la fin de mat�riel m�tal nu, les Javascripts ex�cutables du navigateur.

Lei Feng r�seau AI Technology Review a appris, le blog de TVM est le premier d�crit le papier:

� Avec l'aide de TVM, les d�veloppeurs ont besoin seulement une petite quantit� de travail suppl�mentaire, vous pouvez facilement courir sur le c�t� du t�l�phone mobile, les appareils embarqu�s et m�me sur la profondeur des t�ches d'apprentissage .TVM navigateur fournit �galement une �tude en profondeur uniforme et la charge de travail sur plusieurs plates-formes mat�rielles cadre d'optimisation, y compris le recours � l'informatique nouveau acc�l�rateur d�di� primitif �.

Et aujourd'hui, Chen Tianqi a publi� une nouvelle dynamique sur le microblogging au didacticiel suivant Tucson Hu Wei a pr�sent� en se concentrant sur la promotion de l'optimisation des op apprentissage en profondeur de TVM.

� Apprentissage en profondeur l'optimisation des op est tr�s importante, mais la question difficile. De Tucson avenir Hu Wei a �crit un tutoriel d�crit comment optimiser l'utilisation de l'�tude approfondie VRA de op gpu, que existante tf passer � travers quelques dizaines de lignes de code pour atteindre vingt-trois python fois plus �lev� ".

Cet article est �galement mis � jour en m�me temps sur le blog TVM, Lei Feng r�seau AI Technology Review premi�re fois � faire la couverture et les rapports.

Hu Wei, ma�trise en g�nie �lectronique, Universit� de Beijing de l'a�ronautique et de l'astronautique, actuellement Gap ann�e, et maintenant la pratique de groupe HPC futur Tucson. L'article, intitul� � Optimiser les GPU profondes d'apprentissage avec les op�rateurs: TVM Un exemple Convolution la profondeur � (� Convolution sens de la profondeur, par exemple, pour optimiser l'utilisation de l'op�rateur GPU TVM apprentissage en profondeur)

op�rateur efficace syst�me d'apprentissage de l'apprentissage en profondeur est la profondeur de base. Ces op�rateurs habituellement difficiles � optimiser, les experts HPC doivent payer beaucoup d'efforts. TVM comme une pile de bout en tenseur IR / DSL, peut faciliter l'ensemble du processus.

Cet article fournit une bonne r�f�rence, comment les d�veloppeurs d'apprendre � �crire noyau GPU haute performance avec l'aide des op�rateurs de TVM. L'�quipe utilise Convolution sens de la profondeur (c.-�-topi.nn.depthwise_conv2d_nchw) � titre d'exemple, et montre comment vous pouvez am�liorer le d�j� optimis� manuellement tensorflow le noyau CUDA.

La description de l'article en utilisant la version finale 2-4 TVM fois plus rapide que tf-1.2 sous diff�rentes charges d'exploitation optimiser le noyau, la fusion de l'op�rateur trois fois plus vite -7 fois. Voici les r�sultats des tests dans le cadre GTX1080, la taille du filtre = , stride = , padding = 'm�me':

Est une circonvolution la profondeur id�e de base de la construction d'un mod�le peut effectivement r�duire la complexit� de calcul de la profondeur des r�seaux de neurones, y compris Xception et MobileNet appartiennent Convolution Google sens de la profondeur.

Dans l'environnement TVM, en cours d'ex�cution du code est Convolution comme la profondeur suivante:

# Rembourrage stagePaddedInput = tvm.compute (

(Lot, in_channel, height_after_pad, width_after_pad),

lambda b, c, i, j: tvm.select (

tvm.all (i > = Pad_top, i - pad_top = pad_left, j - Entr�e pad_left , Tvm.const (0,0)),

name = "PaddedInput") # depthconv stagedi = tvm.reduce_axis ((0, filter_height), name = 'di') dj = tvm.reduce_axis ((0, filter_width), name = 'dj') Output = tvm.compute (

(Lot, out_channel, out_height, out_width),

lambda b, c, i, j: tvm.sum (

PaddedInput �* Filtre ,

axe = ),

name = 'DepthwiseConv2d')

Guide GPU g�n�ral Optimisation

Hu Wei a mentionn� trois probl�mes majeurs lors de l'optimisation du code CUDA devrait g�n�ralement �tre not� dans l'article, qui est, la r�utilisation des donn�es (r�utilisation des donn�es), la m�moire partag�e (m�moire partag�e) et une violation d'acc�s (conflits bancaires).

Dans l'architecture de l'informatique moderne, le co�t est calcul� � partir du chargement des donn�es de la m�moire est beaucoup plus �lev� que le simple virgule flottante. Par cons�quent, nous voulons �tre charg� dans un registre ou cache peut �tre utilis� � nouveau dans les donn�es d'entr�e.

Il existe deux formes de la convolution de la profondeur de la r�utilisation des donn�es: des filtres d'entr�e r�utilisation et le recyclage, les anciennes lames sur un canal d'entr�e et calcule une pluralit� de fois, ce qui se produit lorsque la tuile, par exemple � convolution de la profondeur 3x3:

Sans tuile, chaque fil et calcule un des �l�ments de donn�es d'entr�e de charge de sortie 3x3. Un total de 16 fils 9x16 charges.

La tuile, chaque fil calcule 2x24x4 sorties des �l�ments de donn�es d'entr�e et des charges. Un total de 4 fils 16x4 charges.

Et la m�moire partag�e violation d'acc�s

tampon de m�moire partag�e peut �tre consid�r�e comme un GPU, et est plus rapide sur la feuille. La pratique habituelle consiste � charger les donn�es de la m�moire globale en m�moire partag�e, et toutes les discussions dans les donn�es de bloc sont lues � partir de la m�moire partag�e.

Afin d'�viter une violation d'acc�s, un fil continu d'un acc�s continu � la meilleure adresse m�moire comme indiqu� (chaque couleur repr�sente une banque de m�moire partag�e):

Pour plus de d�tails se r�f�rer https://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/

processus d'optimisation sp�cifique

Calcul de l'entr�e ligne de remplissage pour enregistrer l'allocation de m�moire

Rembourrage est explicitement d�clar� comme une �tape distincte. En calculant l'allocation de m�moire en ligne pour �viter la redondance:

s = tvm.create_schedule (s) Output.op .compute_inline

Le passage d'un grand bloc en plus petits

Une approche simple est un traitement de CUDA un bloc de canal d'entr�e et les filtres correspondants, pour calculer la m�moire partag�e apr�s le chargement:

IS = s.cache_read (PaddedInput, "partag�e", )

FS = s.cache_read (Filter, "partag�e", )

block_y = tvm.thread_axis ( "blockIdx.y")

block_x = tvm.thread_axis ( "blockIdx.x")

# Bind la dimension du lot (N dans NCHW) avec block_y

s .bind (Output.op.axis , block_y)

# Bind la dimension du canal (C en NCHW) avec block_x

s .bind (Output.op.axis , block_x)

La figure ci-dessous montre les r�sultats des tests, le co�t du temps moyen de fonctionnement des GTX 10801000 fois et par rapport � tensorflow et depthwise_conv2d.

la performance est bonne, mais si elle est de 64 x 64, alors la performance sera grandement diminu�e si le canal est de 21 x 21 ou 32 x 32 taille,. Si vous faites des changements, l'effet augmentera beaucoup:

Le nombre de fils param�tres de r�glage

Enfilez r�alis� dans un cuda bloc 32 x 32, comme suit:

Comment num_thread_y et num_thread_x ces deux param�tres ajust�s pour obtenir la solution optimale? Dans Filter = et foul�e = ci-dessous:

Gr�ce � des tests, l'�quipe a obtenu les r�sultats suivants:

� grande �chelle la r�utilisation des donn�es de tuiles est bonne, mais pas propice � la lecture de la m�moire locale.
Diff�rents effets num_thread_y et l'acc�s num_thread_x au conflit.
Et num_thread_x num_thread_y meilleure combinaison d'acc�s � la m�moire partag�e n�cessaire pour atteindre (zone de stockage, �viter des conflits) efficaces, la r�utilisation des donn�es, et l'�quilibre de la m�moire locale lue.

Par la recherche de la force brute, la TVM nous pouvons num_thread_y et num_thread_x pass�s comme arguments � l'annexe de la fonction, et d'essayer toutes les combinaisons possibles pour trouver la combinaison optimale.

Vthread (fil virtuel) et Mod�les strided

En TVM, Vthread peut soutenir efficacement mod�les strided.

Dans le filtre en cas = , foul�e = , blocking_h = 32, blocking_w = 32, les r�sultats sont comme suit:

plus vite que le cas 1 cas 2, �tant donn� que le bo�tier 2 num_thread_x = 8 et num_vthread_x = 4 cas, afin d'assurer des fils continus acc�der � des adresses de m�moire cons�cutives, pour des conflits d'acc�s � �viter, comme le montre (chaque couleur repr�sente un fil charge de travail):

Rappelons � nouveau le contraste et tensorflow:

fusion des op�rateurs

op�rateur de fusion est un proc�d� typique pour l'optimisation de r�seau d'apprentissage de la profondeur, dans la TVM en tenant compte du mod�le original depthwise_conv2d + + scale_shift de Relu, peut �tre l�g�rement modifi� comme suit:

IR g�n�r� comme suit:

/ * Entr�e = , Filter = , foul�e = , padding = 'SAME' * / {produire Relu

// attr �thread_extent = 1 // attr �storage_scope = allouer DepthwiseConv2d "local"

// attr �thread_extent = 1 // attr �thread_extent = 8 // attr �thread_extent = 8 produits DepthwiseConv2d {

pour (i, 0, 4) {

pour (j, 0, 4) {

DepthwiseConv2d = 0.000000f

pour (di, 0, 3) {

pour (dj, 0, 3) {

DepthwiseConv2d = (DepthwiseConv2d + (tvm_if_then_else (((((((1 - di) - i)}

}

pour (i2.inner.inner.inner, 0, 4) {

pour (i3.inner.inner.inner, 0, 4) {

Relu = max (((DepthwiseConv2d * Echelle ) + Maj ), 0.000000f)

}

}}

Vous pouvez le voir, chaque fil avant que les r�sultats sont �crits depthwise_conv2d m�moire globale, et calcule scale_shift Relu. Fusion avec un seul op�rateur depthwise_conv2d aussi vite. Ci-dessous = entr�e , Filter = , stride = , les r�sultats rembourrage = 'SAME' est:

tf-1.2 depthwise_conv2d: nous 251,6
tf-1,2 depthwise_conv2d + scale_shift + Relu (s�par�e): 419,9 nous
TVM depthwise_conv2d: nous 90,9
TVM depthwise_conv2d + scale_shift + Relu (fusion): 91,5 nous

code plus optimis� peut se r�f�rer aux liens suivants:

Declare: https://github.com/dmlc/tvm/blob/master/topi/python/topi/nn/convolution.py

Horaire: https://github.com/dmlc/tvm/blob/master/topi/python/topi/cuda/depthwise_conv2d.py

Test: https://github.com/dmlc/tvm/blob/master/topi/recipe/conv/depthwise_conv2d_test.py