C++, Qt, OpenGL, CUDA

Cette page est en lecture seule. Vous pouvez afficher le texte source, mais pas le modifier. Contactez votre administrateur si vous pensez qu'il s'agit d'une erreur.

====== Optimiser les transferts de données ======

Optimisation de la mémoire utilisée (pinned memory), l'overlapping (copie asynchrone), zero-copy, Unified Virtual Addressing

===== Pinned memory =====

Pageable memory
: mémoire dans l'hôte, qui est susceptible d'être copiée sur le disque dur (mémoire virtuelle) lorsque la RAM est saturée. Permet d'utiliser plus de mémoire qu'il n'est possible qu'avec la RAM.


Pinned memory
: mémoire sur l'hôte qui ne peut pas être copiée sur le disque dur.

Demande d'accès GPU : par DMA (direct memory access), page-lock de la mémoire (interdit la copie), copie temporaire dans mémoire pinned (pinned host buffer). Pour éviter le problème, utilisant directement la pinned memory (taille plus importante que le pinned buffer) avec cudaHostAlloc (réservation et allocation) ou cudaHostRegister (mémoire déjà allouée).

Vérifier le type de mémoire utilisée : dans Command Line Profiler avec l'option memtransferhostmemtype. 0 = pageable memory, 1 = pinned memory.

Peu de pinned memory disponible, à utiliser avec précaution (peut ralentir le système complet). Quelle est la quantité de pinned memory disponible ?

===== Regouper les petits transferts de données =====

Surcoût fixe de transfert de données (latence). Regroupement des blocks de données permet de limiter cet overhead. Tester plusieurs taille de transfert.

{{ :460-1_taille_blocs_memoire.png |}}

<note legende>Bande passante en fonction de la taille des blocs mémoire</note>

===== Transferts asynchrones =====

Copie mémoire normalement bloquant. Possibilité d'utiliser non bloquant

  * cudaMemcpy() : synchrone, utilise pageable ou pinned memories
  * cudaMemcpyAsync() : asynchrone, utilise pinned memory et streams


interleave : croise les transferts de données. overlapping : masque les transferts mémoires (réalisés en même temps que les calculs). Overlapping transfert et calculs possible sur toutes les versions de CUDA. Version récentes de CUDA : overlapping transferts de données in et out et calculs. Sur streams différents

Comment tester si overlapping possible ? (et quel type d'ovelapping ?) : cudaDeviceProp.deviceOverlap.

Remarque : stream par défaut (0) est bloquant pour tous les streams (il attend que tous les autres streams ont fini avant de travailler)

Voir avec le profiler comment fonctionne ces streams et l'overlapping :

{{ :460-2_sans_stream_.png |}}

<note legende>Transfert sans overlapping</note>

{{ :460-3_interleave.png |}}

<note legende>Transfert avec interleave</note>

{{ :460-4_overlapping_simple.png |}}

{{ :460-5_overlapping_simple.png |}}

<note legende>Overlapping simple (mémoires et caluls)</note>

{{ :460-6_overlapping_double.png |}}

{{ :460-7_overlapping_double.png |}}

<note legende>Overlapping double</note>

Combien de streams utilisables ?
 
Lancement des instructions : lancer tous les commandes par streams (copy host-device, calculs, copy device-host) ou lancer par commandes (copy host-device sur tous les streams, puis tous les calculs sur tous les streams, puis tous les copy device-host sur tous les streams).

===== Burst memory =====

Accès aux banks de la mémoire globale : par paquet de 128 bits (à vérifier) = transaction. Sans burst :

{{ :460-8_no_burst.png |}}

<note legende>Sans burst</note>

Avec burst, si données coalescente :

{{ :460-9_burst_simple.png |}}

<note legende>Burst simple</note>

Utilisation de plusieurs banks pour remplir entre les burst :

{{ :460-10_multi-bank_burst.png |}}

<note legende>Multi-bank burst</note>

Avec les multi-bank :

{{ :460-11_multi-bank.png |}}

<note legende>Multi-bank et memory coalescence</note>

Problèmes : non alignement des données sur les bursts. Nécessite plusieurs requêtes DRAM au lieu de 1. Mesure efficacité : numbre en pratique / nb théorique de burst. Coalscence si :

<code cpp>
A[(terms independent of threadIdx.x)+ threadIdx.x];
</code>

{{ :460-12_align_et_non_aligne.png |}}

<note legende>Alignement</note>

Par exemple, transformation 2D -> 1D. Accès dans A non coalescent, accès dans B coalescent. Utilisation de sharde memory.

image ???

<note legende>Conversion matrice 2D en mémoire 1D</note>
C++, Qt, OpenGL, CUDA

Outils d'utilisateurs

Outils du Site

Outils de la Page
