integer - number - 256 bytes

¿Número entero de 128 bits en cuda?

Question

Me las arreglé para instalar mi cuda SDK bajo Linux Ubuntu 10.04. Mi tarjeta gráfica es una NVIDIA geForce GT 425M, y me gustaría usarla para algunos problemas computacionales pesados. Lo que me pregunto es: ¿hay alguna forma de utilizar algunos int bits de 128 bits sin signo? Cuando uso gcc para ejecutar mi programa en la CPU, estaba usando el tipo __uint128_t, pero usarlo con cuda no parece funcionar. ¿Hay algo que pueda hacer para tener enteros de 128 bits en cuda?

Muchas gracias Matteo Monti Msoft Programming

Answer 1

CUDA no admite enteros de 128 bits de forma nativa. Puede simular las operaciones usted mismo utilizando dos enteros de 64 bits.

Mira esta publicación :

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

Answer 2

Para obtener el mejor rendimiento, uno querría mapear el tipo de 128 bits sobre un tipo de vector CUDA adecuado, como uint4, e implementar la funcionalidad mediante el ensamblaje PTX en línea. La adición se vería así:

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 multiplicación se puede construir de manera similar utilizando el ensamblaje PTX en línea dividiendo los números de 128 bits en fragmentos de 32 bits, calculando los productos parciales de 64 bits y agregándolos de manera apropiada. Obviamente esto toma un poco de trabajo. Uno podría obtener un rendimiento razonable en el nivel C al dividir el número en fragmentos de 64 bits y usar __umul64hi () junto con la multiplicación regular de 64 bits y algunas adiciones. Esto daría como resultado lo siguiente:

__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;
}

A continuación se muestra una versión de la multiplicación de 128 bits que utiliza el ensamblaje PTX en línea. Requiere PTX 3.0, que se envió con CUDA 4.2, y el código requiere una GPU con al menos capacidad de cálculo 2.0, es decir, un dispositivo de clase Fermi o Kepler. El código usa la cantidad mínima de instrucciones, ya que se necesitan dieciséis multiplicaciones de 32 bits para implementar una multiplicación de 128 bits. En comparación, la variante anterior que utiliza CUDA intrínseca recopila 23 instrucciones para un objetivo 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"(a.x), "r"(a.y), "r"(a.z), "r"(a.w),
           "r"(b.x), "r"(b.y), "r"(b.z), "r"(b.w));
    return res;
}