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?
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