Numero intero a 128 bit su cuda?

Sono appena riuscito a installare il mio cuda SDK sotto Linux Ubuntu 10.04. La mia scheda grafica è una NVIDIA geForce GT 425M, e mi piacerebbe usarla per qualche problema computazionale. Quello che mi chiedo è: c’è un modo per usare un int var a 128 bit senza segno? Usando gcc per eseguire il mio programma sulla CPU, stavo usando il tipo __uint128_t, ma il suo utilizzo con cuda non sembra funzionare. C’è qualcosa che posso fare per avere interi a 128 bit su cuda?

Grazie mille Matteo Monti Msoft Programming

Per prestazioni ottimali, si consiglia di mappare il tipo a 128 bit su un tipo di vettore CUDA adatto, come uint4, e implementare la funzionalità utilizzando l’assembly inline PTX. L’aggiunta sarebbe simile a questa:

typedef uint4 my_uint128_t; __device__ my_uint128_t add_uint128 (my_uint128_t addend, my_uint128_t augend) { my_uint128_t res; asm ("add.cc.u32 %0, %4, %8;\n\t" "addc.cc.u32 %1, %5, %9;\n\t" "addc.cc.u32 %2, %6, %10;\n\t" "addc.u32 %3, %7, %11;\n\t" : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w) : "r"(addend.x), "r"(addend.y), "r"(addend.z), "r"(addend.w), "r"(augend.x), "r"(augend.y), "r"(augend.z), "r"(augend.w)); return res; } 

La moltiplicazione può essere similmente costruita utilizzando l’assembly inline PTX suddividendo i numeri a 128 bit in blocchi a 32 bit, calcolando i prodotti parziali a 64 bit e aggiungendoli in modo appropriato. Ovviamente questo richiede un po ‘di lavoro. È ansible ottenere prestazioni ragionevoli a livello C suddividendo il numero in blocchi di 64 bit e utilizzando __umul64hi () in combinazione con la normale moltiplicazione a 64 bit e alcune aggiunte. Ciò comporterebbe quanto segue:

 __device__ my_uint128_t mul_uint128 (my_uint128_t multiplicand, my_uint128_t multiplier) { my_uint128_t res; unsigned long long ahi, alo, bhi, blo, phi, plo; alo = ((unsigned long long)multiplicand.y << 32) | multiplicand.x; ahi = ((unsigned long long)multiplicand.w << 32) | multiplicand.z; blo = ((unsigned long long)multiplier.y << 32) | multiplier.x; bhi = ((unsigned long long)multiplier.w << 32) | multiplier.z; plo = alo * blo; phi = __umul64hi (alo, blo) + alo * bhi + ahi * blo; res.x = (unsigned int)(plo & 0xffffffff); res.y = (unsigned int)(plo >> 32); res.z = (unsigned int)(phi & 0xffffffff); res.w = (unsigned int)(phi >> 32); return res; } 

Di seguito è riportata una versione della moltiplicazione a 128 bit che utilizza l’assembly inline PTX. Richiede PTX 3.0, fornito con CUDA 4.2, e il codice richiede una GPU con almeno la capacità di calcolo 2.0, cioè un dispositivo di class Fermi o Kepler. Il codice utilizza il numero minimo di istruzioni, in quanto sono necessari sedici multipli a 32 bit per implementare una moltiplicazione a 128 bit. Per fare un confronto, la variante sopra usando le intrinseche CUDA compila a 23 istruzioni per un objective sm_20.

 __device__ my_uint128_t mul_uint128 (my_uint128_t a, my_uint128_t b) { my_uint128_t res; asm ("{\n\t" "mul.lo.u32 %0, %4, %8; \n\t" "mul.hi.u32 %1, %4, %8; \n\t" "mad.lo.cc.u32 %1, %4, %9, %1;\n\t" "madc.hi.u32 %2, %4, %9, 0;\n\t" "mad.lo.cc.u32 %1, %5, %8, %1;\n\t" "madc.hi.cc.u32 %2, %5, %8, %2;\n\t" "madc.hi.u32 %3, %4,%10, 0;\n\t" "mad.lo.cc.u32 %2, %4,%10, %2;\n\t" "madc.hi.u32 %3, %5, %9, %3;\n\t" "mad.lo.cc.u32 %2, %5, %9, %2;\n\t" "madc.hi.u32 %3, %6, %8, %3;\n\t" "mad.lo.cc.u32 %2, %6, %8, %2;\n\t" "madc.lo.u32 %3, %4,%11, %3;\n\t" "mad.lo.u32 %3, %5,%10, %3;\n\t" "mad.lo.u32 %3, %6, %9, %3;\n\t" "mad.lo.u32 %3, %7, %8, %3;\n\t" "}" : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w) : "r"(ax), "r"(ay), "r"(az), "r"(aw), "r"(bx), "r"(by), "r"(bz), "r"(bw)); return res; } 

CUDA non supporta gli interi a 128 bit in modo nativo. Puoi falsificare le operazioni manualmente usando due interi a 64 bit.

Guarda questo post :

 typedef struct { unsigned long long int lo; unsigned long long int hi; } my_uint128; my_uint128 add_uint128 (my_uint128 a, my_uint128 b) { my_uint128 res; res.lo = a.lo + b.lo; res.hi = a.hi + b.hi + (res.lo < a.lo); return res; } 

Una risposta molto ritardata, ma potresti prendere in considerazione l’utilizzo di questa libreria:

https://github.com/curtisseizert/CUDA-uint128

che definisce una struttura di dimensioni di 128 bit, con metodi e funzioni di utilità indipendenti per farlo funzionare come previsto, che consente di utilizzarlo come un normale intero. Soprattutto.