2012-09-27 18 views
28

Perché lo standard atomicAdd() per i doppi è stato implementato esplicitamente come parte di CUDA 4.0 o superiore?Perché atomicAdd non è stato implementato per i doppi?

Dall'appendice F Pagina 97 del CUDA programming guide 4.1 sono state implementate le seguenti versioni di atomicAdd.

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); 
float atomicAdd(float* address, float val) 

La stessa pagina va avanti per dare una piccola implementazione di atomicAdd per il doppio come segue che ho appena iniziato a utilizzare nel mio progetto.

__device__ double atomicAdd(double* address, double val) 
{ 
    unsigned long long int* address_as_ull = 
          (unsigned long long int*)address; 
    unsigned long long int old = *address_as_ull, assumed; 
    do { 
     assumed = old; 
old = atomicCAS(address_as_ull, assumed, 
         __double_as_longlong(val + 
           __longlong_as_double(assumed))); 
    } while (assumed != old); 
    return __longlong_as_double(old); 
} 

Perché non definire il codice sopra come parte di CUDA?

+2

Probabilmente in modo che ogni utente ne sia consapevole dell'implementazione, in quanto non è un'istruzione incorporata e la logica dei tentativi può essere soggetta a livelock (poiché non vi è alcuna garanzia di equità, un thread può essere bloccato per finché ci sono altri thread che aggiornano la stessa variabile). – tera

risposta

31

Modifica: A partire da CUDA 8, la doppia precisione atomicAdd() è implementata in CUDA con supporto hardware nelle GPU SM_6X (Pascal).

Attualmente, nessun dispositivo CUDA supporta atomicAdd per double in hardware. Come si è notato, può essere implementato in termini di atomicCAS su numeri interi a 64 bit, ma c'è un costo di prestazione non banale per questo.

Pertanto, il team del software CUDA ha scelto di documentare un'implementazione corretta come opzione per gli sviluppatori, piuttosto che renderla parte della libreria standard CUDA. In questo modo gli sviluppatori non stanno inconsapevolmente optando per un costo di performance che non capiscono.

A parte: non penso che questa domanda debba essere chiusa come "non costruttiva". Penso che sia una domanda perfettamente valida, +1.

+1

Sì, ma tecnicamente sei una delle poche persone in grado di rispondere alla domanda. Mentre ho detto perché penso che abbia molto senso in questo modo, solo tu puoi dire se questo è il motivo per cui il team CUDA l'ha scelto in questo modo. ;-) Ad ogni modo non ero io a sviare la domanda. – tera

+0

Ci sono più persone NVIDIA che leggono e rispondono alle domande CUDA su SO (specialmente mentre i nostri forum degli sviluppatori non funzionano), e questo fatto rende valide domande come questa. E avresti potuto pubblicare il tuo commento come risposta, e sarebbe stato corretto, e lo avrei svalutato. :) A proposito, non presumo che hai downvoted; Mi riferivo all'unico voto per chiudere la domanda. – harrism

+1

Sono d'accordo, questa è una domanda perfettamente valida, le intestazioni CUDA potrebbero aver implementato la doppia atomica nel software. Anche se il modo in cui è stato formulato ha innescato la luce rossa per alcune persone, penso che la decisione dovrebbe essere annullata! – pszilard

Problemi correlati