Emulazione di FP64 con 2 FP32 su una GPU

Question

Se si dovesse emulare un virgola mobile a doppia precisione con due punti di virgola mobile di precisione singoli, quale sarebbe la prestazione e può essere eseguita correttamente?

Attualmente Nvidia si sta caricando piuttosto bene per le schede Tesla abilitate a doppia precisione che consentono di ottenere un terzo delle prestazioni di precisione singole (eccezioni notevoli Titan/Titan Black).

Se si dovesse utilizzare una GPU Geforce con doppia precisione snodata ed emulare una precisione doppia usando 2 galleggianti di precisione singoli come sarebbe la prestazione?

Answer 1

È possibile ottenere una stima approssimativa della prestazione contando il numero di operazioni float necessarie per implementare ogni operazione a doppio flottante. Si consiglia di ispezionare il codice binario con cuobjdump --dump-sass per ottenere un conteggio accurato. Sto mostrando una moltiplicazione a doppio float in basso che sfrutta appieno il supporto FMA (fuse multiply-add) sulla GPU. Per il codice di aggiunta a virgola mobile, vorrei indirizzarti a a paper by Andrew Thall perché non ho il tempo di farlo ora. Dall'analisi precedente ritengo che il codice di addizione fornito nel documento sia corretto e che eviti le insidie ​​più comuni in implementazioni più veloci ma meno accurate (che perdono la precisione quando la grandezza degli operandi è entro un fattore due).

Se sei uno sviluppatore registrato CUDA è possibile scaricare il codice doppia-doppia dal sito degli sviluppatori di NVIDIA (log in a https://developer.nvidia.com) che è sotto licenza BSD, e rielaborare relativamente rapidamente in codice doppio float. Il doppio doppio codice di NVIDIA supporta le operazioni di addizione, sottrazione, divisione, radice quadrata e radice quadrata reciproca.

Come si può vedere, la moltiplicazione sotto richiede 8 float istruzioni; la negazione unaria è assorbita in FMA. L'aggiunta richiede circa 20 float istruzioni. Tuttavia, le sequenze di istruzioni per le operazioni a doppio movimento richiedono anche variabili temporanee, che aumentano la pressione del registro e possono ridurre l'occupazione. Una stima ragionevolmente conservativa può quindi essere che l'aritmetica a doppio galleggiamento esegua a 1/20 il throughput dell'aritmetica nativa float. Puoi facilmente misurarlo tu stesso, nel contesto pertinente per te, ad esempio il tuo caso/i d'uso.

typedef float2 dblfloat; // .y = head, .x = tail 

__host__ __device__ __forceinline__ 
dblfloat mul_dblfloat (dblfloat x, dblfloat y) 
{ 
    dblfloat t, z; 
    float sum; 
    t.y = x.y * y.y; 
    t.x = fmaf (x.y, y.y, -t.y); 
    t.x = fmaf (x.x, y.x, t.x); 
    t.x = fmaf (x.y, y.x, t.x); 
    t.x = fmaf (x.x, y.y, t.x); 
    /* normalize result */ 
    sum = t.y + t.x; 
    z.x = (t.y - sum) + t.x; 
    z.y = sum; 
    return z; 
} 

Si noti che in varie applicazioni, l'aritmetica con doppio float completo potrebbe non essere necessaria. Invece si può usare il calcolo float, aumentato dalle tecniche di compensazione degli errori, uno dei più vecchi dei quali è il Kahan summation. Ho dato una breve panoramica della letteratura facilmente disponibile su tali metodi in un recent posting in the NVIDIA developer forums. Nei commenti sopra, Robert Crovella ha anche indicato uno GTC 2015 talk by Scott LeGrand, che non ho ancora avuto il tempo di verificare.

Per quanto riguarda la precisione, il doppio movimento ha una precisione rappresentazionale di 49 (24 + 24 + 1) bit, rispetto a IEEE-755 double che fornisce 53 bit. Tuttavia, il doppio float non può mantenere questa precisione per gli operandi di piccole dimensioni, poiché la porzione di coda può diventare un denormale o zero. Quando il supporto denormale è attivato, i 49 bit di precisione sono garantiti per 2 ^-101 < = | x | . Il supporto Denormal per float è attivato per impostazione predefinita nella catena di strumenti CUDA per architetture> = sm_20, il che significa che tutte le architetture supportate dalla versione attualmente in commercio, CUDA 7.0.

Diversamente dalle operazioni sui dati IEEE-754 double, le operazioni a doppio movimento non sono arrotondate correttamente.Per la moltiplicazione a doppio float sopra, utilizzando 2 miliardi di casi di test casuali (con tutti gli operandi sorgente e i risultati entro i limiti sopra indicati), ho osservato un limite superiore di 1.42e-14 per l'errore relativo. Non ho dati per l'aggiunta a doppio float, ma il suo limite di errore dovrebbe essere simile.