entero de 128 bits en cuda?

Question

Acabo de 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 Programación

Answer 1

Para un mejor rendimiento, que uno quiere representar el tipo de 128 bits en la parte superior de un CUDA tipo vector adecuado, tal como uint4, e implementar la funcionalidad mediante el montaje PTX en línea . La adición sería algo como esto:

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 de manera similar puede construirse utilizando ensamblador en línea PTX rompiendo los números de 128 bits en trozos de 32 bits, el cálculo de los productos parciales de 64 bits y añadiendo 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 el 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 ensamblado en línea PTX. 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; 
} 

Answer 2

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

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