cuda download - aritmetica modulare sulla gpu




nvidia toolkit (4)

Qualche tempo fa ho sperimentato molto con l'aritmetica modulare sulla GPU. Sulle GPU Fermi è possibile utilizzare l'aritmetica a doppia precisione per evitare costose operazioni di div e mod. Ad esempio, la moltiplicazione modulare può essere eseguita come segue:

// fast truncation of double-precision to integers
#define CUMP_D2I_TRUNC (double)(3ll << 51)
// computes r = a + b subop c unsigned using extended precision
#define VADDx(r, a, b, c, subop) \
    asm volatile("vadd.u32.u32.u32." subop " %0, %1, %2, %3;" :  \
            "=r"(r) : "r"(a) , "r"(b), "r"(c));

// computes a * b mod m; invk = (double)(1<<30) / m
__device__ __forceinline__ 
unsigned mul_m(unsigned a, unsigned b, volatile unsigned m,
    volatile double invk) { 

   unsigned hi = __umulhi(a*2, b*2); // 3 flops
   // 2 double instructions
   double rf = __uint2double_rn(hi) * invk + CUMP_D2I_TRUNC;
   unsigned r = (unsigned)__double2loint(rf);
   r = a * b - r * m; // 2 flops

   // can also be replaced by: VADDx(r, r, m, r, "min") // == umin(r, r + m);
   if((int)r < 0) 
      r += m;
   return r;
}

Tuttavia questo funziona solo per i moduli interi a 31 bit (se 1 bit non è fondamentale per te) e devi anche precomputare preventivamente 'invk'. Questo dà il minimo assoluto di istruzioni che posso raggiungere, es .:

SHL.W R2, R4, 0x1;
SHL.W R8, R6, 0x1;
IMUL.U32.U32 R4, R4, R6;
IMUL.U32.U32.HI R8, R2, R8;
I2F.F64.U32 R8, R8;
DFMA R2, R2, R8, R10;
IMAD.U32.U32 R4, -R12, R2, R4;
ISETP.GE.AND P0, pt, R4, RZ, pt;
@!P0 IADD R4, R12, R4;

Per la descrizione dell'algoritmo, puoi dare un'occhiata al mio articolo: gpu_resultants . Altre operazioni come (x y - z w) mod m sono anche spiegate qui.

Per curiosità, ho confrontato le prestazioni dell'algoritmo risultante usando la moltiplicazione modulare:

unsigned r = (unsigned)(((u64)a * (u64)b) % m);

contro la versione ottimizzata con mul_m.

Aritmetica modulare con operazione% predefinita:

low_deg: 11; high_deg: 2481; bits: 10227
nmods: 330; n_real_pts: 2482; npts: 2495

res time: 5755.357910 ms; mod_inv time: 0.907008 ms; interp time: 856.015015 ms; CRA time: 44.065857 ms
GPU time elapsed: 6659.405273 ms; 

Aritmetica modulare con mul_m:

low_deg: 11; high_deg: 2481; bits: 10227
nmods: 330; n_real_pts: 2482; npts: 2495

res time: 1100.124756 ms; mod_inv time: 0.192608 ms; interp time: 220.615143 ms; CRA time: 10.376352 ms
GPU time elapsed: 1334.742310 ms; 

Quindi, in media, è circa 5 volte più veloce. Nota anche che potresti non vedere un'accelerazione se hai appena valutato le prestazioni aritmetiche crude usando un kernel con una serie di operazioni mul_mod (come l' esempio di saxpy ). Ma nelle applicazioni reali con logica di controllo, barriere di sincronizzazione, ecc., L'accelerazione è molto evidente.

Sto lavorando sull'algoritmo della GPU che dovrebbe fare molti calcoli modulari. In particolare, varie operazioni su matrici in un campo finito che a lungo termine riducono a operazioni primitive come: (a * b - c * d) mod m o (a * b + c) mod m dove a, b, c d sono residui modulo m ed m è un primo di 32 bit.

Attraverso la sperimentazione ho appreso che le prestazioni dell'algoritmo sono per lo più limitate dall'aritmetica modulare lenta perché il modulo intero (%) e le operazioni di divisione non sono supportate sulla GPU nell'hardware.

Apprezzo se qualcuno mi può dare un'idea di come realizzare computazioni modulari efficienti con CUDA?

Per vedere come questo è implementato su CUDA, io uso il seguente frammento di codice:

__global__ void mod_kernel(unsigned *gout, const unsigned *gin) {

unsigned tid = threadIdx.x;
unsigned a = gin[tid], b = gin[tid * 2], m = gin[tid * 3];

typedef unsigned long long u64;

__syncthreads();
unsigned r = (unsigned)(((u64)a * (u64)b) % m);
__syncthreads();
gout[tid] = r;
}

Questo codice non dovrebbe funzionare, volevo solo vedere come la riduzione modulare è implementata su CUDA.

Quando smonto questo con cuobjdump --dump-sass (grazie njuffa per un consiglio!), Vedo quanto segue:

/*0098*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
/*00a0*/     /*0x1c315c4350000000*/     IMUL.U32.U32.HI R5, R3, R7;
/*00a8*/     /*0x1c311c0350000000*/     IMUL.U32.U32 R4, R3, R7;
/*00b0*/     /*0xfc01dde428000000*/     MOV R7, RZ;
/*00b8*/     /*0xe001000750000000*/     CAL 0xf8;
/*00c0*/     /*0x00000007d0000000*/     BPT.DRAIN 0x0;
/*00c8*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;

Si noti che tra le due chiamate a bar.red.popc c'è una chiamata alla procedura 0xf8 che implementa qualche sofisticato algoritmo (circa 50 istruzioni o anche più). Non sorprende che l'operazione mod (%) sia lenta


Una GPU Fermi di fascia alta (ad esempio una GTX 580) probabilmente ti darà le migliori prestazioni tra le carte di spedizione per questo. Vorresti che tutti gli operandi a 32 bit fossero di tipo "unsigned int" per le migliori prestazioni, in quanto vi è un ulteriore sovraccarico per la gestione delle divisioni e dei moduli firmati.

Il compilatore genera un codice molto efficiente per la divisione e modulo con divisore fisso. Ricordo che di solito sono circa tre o cinque istruzioni di istruzioni della macchina su Fermi e Kepler. È possibile controllare il SASS generato (codice macchina) con cuobjdump --dump-sass. Potresti essere in grado di utilizzare funzioni basate su modelli con divisori costanti se usi solo pochi divisori diversi.

Dovresti vedere nell'ordine di sedici istruzioni SASS in linea generate per le operazioni a 32 bit senza segno con divisore variabile, attraverso Fermi e Kepler. Il codice è limitato dal throughput di numeri interi e le GPU di classe Fermi sono competitive con le soluzioni hardware. Le prestazioni un po 'ridotte si riscontrano sulle GPU di classe Kepler attualmente in vendita a causa del loro ridotto throughput multiplo intero.

[Aggiunto dopo, dopo aver chiarito la domanda:]

La divisione a 64 bit senza segno e modulo con divisore variabile d'altra parte sono chiamate subroutine di circa 65 istruzioni su Fermi e Kepler. Sembrano vicini a quelli ottimali. Su Fermi, questo è ancora ragionevolmente competitivo con le implementazioni hardware (si noti che le divisioni in interi a 64 bit non sono esattamente super veloci sulle CPU che forniscono questo come un'istruzione integrata). Di seguito è riportato un codice che ho postato nei forum NVIDIA qualche volta indietro per il tipo di attività descritta nel chiarimento. Evita la divisione costosa, ma presuppone che lotti piuttosto ampi di operandi condividano lo stesso divisore. Usa l'aritmetica a doppia precisione, che è particolarmente veloce sulle GPU di classe Tesla (al contrario delle carte di consumo). Ho fatto solo un test rapido del codice, potresti testarlo con più attenzione prima di distribuirlo.

// Let b, p, and A[i] be integers < 2^51
// Let N be a integer on the order of 10000
// for i from 1 to N
// A[i] <-- A[i] * b mod p

/*---- kernel arguments ----*/
unsigned long long *A;
double b, p; /* convert from unsigned long long to double before passing to kernel */
double oop;  /* pass precomputed 1.0/p to kernel */

/*---- code inside kernel -----*/
double a, q, h, l, rem;
const double int_cvt_magic = 6755399441055744.0; /* 2^52+2^51 */

a = (double)A[i];

/* approximate quotient and round it to the nearest integer */
q = __fma_rn (a * b, oop, int_cvt_magic);
q = q - int_cvt_magic;

/* back-multiply, representing p*q as a double-double h:l exactly */
h = p * q;
l = __fma_rn (p, q, -h);

/* remainder is double-width product a*b minus double-double h:l */
rem = __fma_rn (a, b, -h);
rem = rem - l;

/* remainder may be negative as quotient rounded; fix if necessary */
if (rem < 0.0) rem += p;

A[i] = (unsigned long long)rem;

Ci sono trucchi per eseguire in modo efficiente le operazioni di mod, ma se m è solo radix 2.

Per esempio, x mod y == x & (y-1), dove y è 2 ^ n. L'esecuzione di operazioni bit a bit è la più veloce.

Altrimenti, probabilmente una tabella di ricerca? Di seguito è riportato un link sulla discussione sull'efficiente implementazione del modulo. Potrebbe essere necessario implementarlo tu stesso per ottenere il massimo da esso.

Calcolo efficiente di mod


Con così tanto potere non sfruttato, non riesco a vedere come sarebbe rimasto inutilizzato per troppo tempo. La domanda è, però, come verrà utilizzata la GPU per questo. CUDA sembra essere una buona ipotesi per ora, ma altre tecnologie stanno emergendo all'orizzonte che potrebbero renderlo più accessibile dallo sviluppatore medio.

Apple ha recentemente annunciato OpenCL che affermano che è molto più di CUDA, ma piuttosto semplice. Non sono sicuro di cosa farne esattamente, ma il gruppo khronos (I ragazzi che lavorano sullo standard OpenGL) stanno lavorando sullo standard OpenCL e sta cercando di renderlo altamente interoperabile con OpenGL. Ciò potrebbe portare a una tecnologia che è più adatta per il normale sviluppo del software.

È un argomento interessante e, per inciso, sto per iniziare la mia tesi di master sul tema del modo migliore per rendere la potenza della GPU disponibile agli sviluppatori medi (se possibile) con CUDA come obiettivo principale.





cuda gpgpu