[Cuda] aritmetica modulare sulla gpu


Answers

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;
Question

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