XII. Fonctions mathématiques▲
Ces fonctions sont des équivalents des fonctions proposées par la librairie standard du C. Cependant, toutes ne le sont pas. Si ces fonctions sont utilisées dans du code pour l'hôte, alors une implémentation CUDA est utilisée. Sinon, c'est l'implémentation de la librairie standard qui est utilisée.
Toutes ces fonctions existent en deux versions : simple et double précision. Si la version double précision est appelée avec un argument de simple précision, c'est la version de simple précision qui sera appelée.
Voici les 68 fonctions supportées :
- x+y
- x*y
- x/y
- 1/x
- 1/sqrtf(x)
- rsqrtf(x)
- sqrtf(x)
- cbrtf(x)
- hypotf(x,y)
- expf(x)
- exp2f(x)
- exp10f(x)
- expm1f(x)
- logf(x)
- log2f(x)
- log10f(x)
- log1pf(x)
- sinf(x)
- cosf(x)
- tanf(x)
- sincosf(x,sptr,cptr)
- asinf(x)
- acosf(x)
- atanf(x)
- atan2f(y,x)
- sinhf(x)
- coshf(x)
- tanhf(x)
- asinhf(x)
- acoshf(x)
- atanhf(x)
- powf(x,y)
- erff(x)
- erfcf(x)
- lgammaf(x)
- tgammaf(x)
- fmaf(x,y,z)
- frexpf(x,exp)
- ldexpf(x,exp)
- scalbnf(x,n)
- scalblnf(x,l)
- logbf(x)
- ilogbf(x)
- fmodf(x,y)
- remainderf(x,y)
- remquof(x,y,iptr)
- modff(x,iptr)
- fdimf(x,y)
- truncf(x)
- roundf(x)
- rintf(x)
- nearbyintf(x)
- ceilf(x)
- floorf(x)
- lrintf(x)
- lroundf(x)
- llrintf(x)
- llroundf(x)
- signbit(x)
- isinf(x)
- isnan(x)
- isfinite(x)
- copysignf(x,y)
- fminf(x,y)
- fmaxf(x,y)
- fabsf(x)
- nanf(cptr)
- nextafterf(x,y)
Toutes ces opérations ne sont pas aussi précises l'une que l'autre : la table des précisions est disponible dans le Programming Guide, dans l'annexe B.
XIII. Fonctions atomiques▲
Ces fonctions ne sont pas disponibles sur tous les périphériques : seulement ceux de Compute Capability supérieure ou égale à 1.1. Ceci se résume ainsi : les cartes GeForce 9000 et supérieures, les Quadro à l'exception des FX 4600 et FX 5600 et les Quadro Plex à l'exception des 1000 et 2100 S4.
Qu'est-ce qu'une fonction atomique ?
Une fonction atomique est une opération ou un ensemble d'opérations qui s'exécute entièrement sans que le thread qui
l'a demandée ne cède sa place pendant son traitement.
Autrement dit, lors d'un appel à une fonction atomique, aucun autre thread CPU ne peut accéder au même processeur
de flux.
Les fonctions atomiques de CUDA se divisent en deux groupes : les fonctions arithmétiques et les fonctions d'opération sur les bits.
Ces fonctions ne peuvent se dérouler qu'en mémoire partagée sur des mots de 32 bits. Seuls les périphériques de Compute Capability de 1.2 ou plus peuvent les effectuer sur des mots de 64 bits. Les périphériques qui en sont capables sont les GeForce GTX 200, les Tesla C1000 et les Quadro FX 5800. Ces périphériques sont aussi les seuls à pouvoir effectuer des calculs à double précision (FP64).
XIII-A. Arithmétiques▲
int atomicAdd(int * address, int val);
unsigned int atomicAdd(unsigned int * address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int * address, unsigned long long int val);Cette fonction lit la variable stockée à l'adresse address et lui ajoute val. La fonction renvoie l'ancienne valeur, avant comparaison.
int atomicSub(int * address, int val);
unsigned int atomicSub(unsigned int * address, unsigned int val);Cette fonction lit la variable stockée à l'adresse address et lui retire val. La fonction renvoie l'ancienne valeur, avant comparaison.
int atomicExch(int * address, int val);
unsigned int atomicExch(unsigned int * address, unsigned int val);
unsigned long long int atomicExch(unsigned long long int * address, unsigned long long int val);
float atomicExch(float * address, float val);Cette fonction lit la variable stockée à l'adresse address et lui met val comme valeur. La fonction renvoie l'ancienne valeur, avant remplacement.
int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address, unsigned int val);Cette fonction lit la variable stockée à l'adresse address et la remplace par le minimum de cette variable et de val. La fonction renvoie l'ancienne valeur, avant calcul.
int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address, unsigned int val);Cette fonction lit la variable stockée à l'adresse address et la remplace par le maximum de cette variable et de val. La fonction renvoie l'ancienne valeur, avant calcul.
unsigned int atomicInc(unsigned int* address, unsigned int val);Cette fonction lit la variable stockée à l'adresse address et l'incrémente si elle est supérieure ou égale à val, sinon l'annule. La fonction renvoie l'ancienne valeur, avant calcul.
unsigned int atomicDec(unsigned int* address, unsigned int val);Cette fonction lit la variable stockée à l'adresse address et la décrémente si elle est supérieure ou égale à val, sinon l'annule. La fonction renvoie l'ancienne valeur, avant calcul.
typedef unsigned int uint;
typedef unsigned long long int ullint,
int atomicCAS(int * address, int compare, int val);
uint atomicCAS(uint * address, uint compare, uint val);
ullint atomicCAS(ullint * address, ullint compare, ullint val);Cette fonction lit la variable stockée à l'adresse address et la remplace par val si elle est égale à compare. La fonction renvoie l'ancienne valeur, avant calcul.
XIII-B. Bits▲
Il n'existe pas une aussi grande quantité d'opérations sur les bits. Seules trois sont supportées par CUDA : et (and), ou (or) et ou exclusif (xor). Le premier paramètre de l'opération est la variable dont un pointeur est passé, le second est la variable passée en argument. La fonction renvoie, comme toujours, la valeur sur le GPU avant calcul.
int atomicAnd(int * address, int val);
unsigned int atomicAnd(unsigned int * address, unsigned int val);
int atomicOr(int * address, int val);
unsigned int atomicOr(unsigned int * address, unsigned int val);
int atomicXor(int * address, int val);
unsigned int atomicXor(unsigned int * address, unsigned int val);XIV. Fonction de synchronisation▲
Dans le runtime que nous étudions, il n'existe qu'une seule fonction de synchronisation des threads : void __syncthreads();.
Elle synchronise tous les threads d'un bloc. Dès qu'ils sont tous arrivés à ce point dans les calculs, ils recommencent.
Cette synchronisation est principalement utilisée pour coordonner la communication entre les threads d'un même bloc. Quand quelques threads accèdent au même emplacement mémoire, il peut y avoir quelques problèmes, vu que l'on ne peut jamais savoir si un thread a déjà écrit ou lu des données. Ceci peut être évité avec les threads : on peut être sûr que tous les threads sont arrivés à cet endroit, qu'ils ont tous lu ou écrit ce qu'ils devaient.
Cette opération est aussi fort coûteuse en temps : elle prend 4 cycles, si elle ne doit attendre aucun thread. Mais elle peut éviter des erreurs d'exécution, qui doivent être évitées à n'importe quel prix.