CUDA Independent Instruction optimization

Question

J'ai besoin de suggestions concernant l'optimisation de mon noyau et du code de l'appareil. Je comprends que la documentation CUDA (et tant de diapositives) suggère l'utilisation de grandes tailles de blocs de thread afin de cacher la mémoire et la latence arithmétique. Mes noyaux et les fonctions de l'appareil sont très intensives en termes de calculs. Par conséquent j'essaie d'utiliser autant de registres que possible et (évidemment) à cause de cela je fais des compromis sur l'occupation. Point est, pour mon application, le parallélisme au niveau de l'instruction est plus important que les gros blocs de thread.

Mais l'idée de base derrière ILP est d'avoir des instructions indépendantes. Ma question est

1) Comment y parvenir? En calcul, il y a toujours des variables qui sont réutilisées pour d'autres calculs.
2) Quelqu'un peut-il suggérer ou fournir des exemples où des instructions dépendantes peuvent être transformées en instructions indépendantes?
3) Je lis aussi (quelque part) que pour le calcul arithmétique, ILP maximum = 4 peut être atteint, c'est-à-dire qu'un thread calcule 4 instructions indépendantes. Est-ce que cela signifie, s'il existe de telles quatre instructions et après cela il y a des instructions dépendantes, warp entrera en attente, jusqu'à ce que les dépendances sont remplies?
4) Quelqu'un peut-il suggérer du matériel de lecture et du code où ILP est exploité? Je présente ici aussi du code pour l'analyse; cela ne veut peut-être rien dire. Le code représente l'équation suivante:

Formula

Le point est que je veux obtenir des performances maximales; et je veux utiliser ILP pour cela. J'ai d'autres fonctions de l'appareil dans mon code; donc je me sers

bloc de fil: 192
14 SM (32 cœurs): 448 (noyaux)
Chaque SM utilise 8 blocs en même temps: 8 x 192: 1536
Lors de la compilation du code avec « -ptxas -options = -v » Je reçois 50 registres par fil (occupation quelque part autour de 33%)

Tous les paramètres utilisés dans l'équation sont de type double (autre que n)
par exemple n = 2. tableau params contient S à param [0] et I1 à param [1] et I2 au param [2]

#define N 3.175e-3 
__device__ double gpu_f_different_mean(double x, double params[], int n) { 

    double S = params[0]; 
    double product_I = 1.0; 

    for (int i = 1; i <= n; i++) { 
     product_I = product_I * params[i]; 
    } 

    double tmp = S * exp(-N * S * x); 
    double outer = product_I * tmp; 

    double result = 0.0; 

    for (int i = 1; i <=n; i++) { 

     double reduction = (params[i] + S * x); 
     double numerator = 1 + N * reduction; 

     double denom_prod = 1.0; 
     for (int j = 1; j<= n; j++) { 
     if (i != j) 
      denom_prod = denom_prod * (params[j] - params[i]); 
     } 

     double denominator = pow(reduction, 2) * denom_prod; 
     result    = result + (numerator/denominator); 
    } 

    return outer * result; 
} 

Hardware

J'utilise Fermi architecture GPU GTX470, calculer la capacité 2,0

Answer 1

plusieurs commentaires:

a) des chaînes de dépendance comme celle provoquée par la mise à jour continue de denom_prod peut être rompu par l'introduction de plusieurs variables de réduction:

double denom_prod1 = 1.0; 
    double denom_prod2 = 1.0; 
    int j; 
    for (j = 1; j <= n-1; j += 2) { 
    if (i != j) 
     denom_prod1 *= (params[j ] - params[i]); 
    if (i != j+1) 
     denom_prod2 *= (params[j+1] - params[i]); 
    } 
    if (j < n) { 
    if (i != j) 
     denom_prod1 = denom_prod * (params[j ] - params[i]); 
    } 
    double denom_prod = denom_prod1 * denom_prod2; 

b) Le conditionnel à l'intérieur de la boucle peut être éliminé en brisant la boucle en deux parties:

double denom_prod = 1.0; 
    for (int j = 1; j < i; j++) 
    denom_prod = denom_prod * (params[j] - params[i]); 
    for (int j = i+1; j <= n; j++) 
    denom_prod = denom_prod * (params[j] - params[i]); 

c) Vous pouvez exploiter le fait que l'échange i et j ne changera pas denom_prod en calculant les résultats pour (i, j) et (j, i) en une fois.

d) reduction * reduction est plus rapide (et potentiellement plus précis) que pow(reduction, 2)

---

En ce qui concerne vos questions:

1) et 2) voir mon commentaire a). 3) Ceci est probablement dû au fait que les GPU de génération Fermi (capacité de calcul 2.x) ont deux ordonnanceurs de distorsion indépendants par SM, chacun capable d'émettre deux instructions par cycle, pour un total de quatre instructions maximum par cycle.

Cependant le problème des instructions dépendantes va plus loin que cela, car les instructions dépendantes souffrent d'une latence de ~ 16,24 cycles. C'est à dire. la seconde des deux instructions dépendantes doit attendre autant de cycles avant de pouvoir être émise. Les cycles intermédiaires peuvent soit être utilisés par des instructions indépendantes provenant de la même chaîne (qui doivent être situées entre les instructions dépendantes, car les GPU Nvidia actuels ne peuvent pas sortir des instructions dans le désordre). Ou ils peuvent être utilisés par des instructions d'autres warps, qui sont toujours indépendantes. Donc, pour une performance optimale, vous voulez soit beaucoup de warps, soit des instructions indépendantes consécutives, ou idéalement les deux.

4) Les publications de Vasily Volkov permettent une excellente lecture sur ce sujet, notamment sa présentation "Better Performance at Lower Occupancy".