Notre précédent article suivait un noyau d'addition vectorielle de nvcc jusqu'aux warps. Cette fois, nous suivons une instruction SASS critique (un chargement global) à travers le matériel d'une RTX4090. Le noyau CUDA étudié comporte deux lignes : __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}. Le SASS compilé montre : /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Ces instructions chargent des éléments du vecteur b depuis la mémoire globale vers les registres.
Une LDG.E demande quatre octets dans chacun des 32 lanes. La servir nécessite quatre secteurs de 32 octets, une ligne de cache, une traduction d'adresse, un passage du crossbar, l'une des trente-six slices L2 et, en cas d'échec partout, une activation et quatre lectures de colonnes sur une puce DRAM. Notre warp vit sur l'une des quatre sous-partitions du SM aux côtés de onze autres warps résidents. L'instruction est émise vers l'unité de chargement/stockage (LSU), qui envoie l'opcode, le masque de lanes actives, les adresses calculées et le numéro de registre au coalescer.
Le coalescer détermine que 4 requêtes de secteurs contiguës sont nécessaires pour les 128 octets demandés par le warp. En entrant dans le cache L1, la requête est envoyée au cache adressé virtuellement, associatif par ensembles 4 voies. La recherche compare les balises dans des ensembles de quatre emplacements. Comme il s'agit d'un premier chargement, il échoue et descend plus loin. Un hit L1 revient en environ 15,4 ns — 40 cycles. La mémoire virtuelle nécessite une traduction avant de quitter le SM. La TLB du SM contient seize traductions récentes. Les échecs coûtent environ 4,4 ns — onze cycles.
Après la traduction, une requête par ligne de 128 octets traverse le crossbar vers l'une des 36 slices L2 de 2 MiB, choisie par une fonction de son adresse physique. Toutes les slices peuvent servir en parallèle, offrant une bande passante agrégée 36 fois supérieure à celle d'une seule slice.
Source: Hacker News · Résumé par HeadlinesBriefing