128-bit-integer auf cuda?

Für die beste Leistung, man möchte Karte die 128-bit-geben Sie auf der Oberseite eines geeigneten CUDA vector Art, wie uint4, und die Funktionalität implementieren, die Verwendung von PTX-inline-Montage. Die Ergänzung würde etwa so Aussehen:

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

Die Multiplikation kann Analog konstruiert werden mit PTX-inline-Montage durch das brechen der 128-bit-zahlen in 32-bit-Blöcken zur Berechnung der 64-bit-teilweise Produkte aus und fügen Sie diese entsprechend. Natürlich braucht ein wenig Arbeit. Könnte man angemessene Leistung bei der C-Ebene durch das brechen der Zahl in 64-bit-Blöcken und mit __umul64hi() in Verbindung mit regelmäßiger 64-bit-Multiplikation und einige Ergänzungen. Dies würde zu folgendem führen:

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

Unten ist eine version der 128-bit-Multiplikation verwendet PTX-inline-Montage. Es erfordert PTX 3.0, ausgeliefert mit CUDA 4.2, und der code erfordert eine GPU mit mindestens compute capability 2.0, d.h. Fermi-oder Kepler-Klasse-Gerät. Der code verwendet die minimale Anzahl von Anweisungen, wie sechzehn 32-bit multipliziert werden benötigt um ein 128-bit-Multiplikation. Durch den Vergleich, die Variante oben mit Hilfe von CUDA-Interna, kompiliert 23 Anweisungen für eine sm_20 Ziel.

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

Eine stark verspätete Antwort, aber könntest du überlegen, die diese Bibliothek nutzen:

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

definiert eine 128-bit-großen Struktur, Methoden und freistehende utility-Funktionen, um es zu funktionieren, wie erwartet, die es erlauben verwendet werden wie eine normale Ganzzahl ist. Meistens.