Для достижения наилучшей производительности желательно отобразить 128-битный тип поверх подходящего векторного типа CUDA, такого как uint4, и реализовать эту функциональность, используя встроенную сборку PTX.Дополнение будет выглядеть примерно так:
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;
}
Аналогично, умножение можно построить с использованием встроенной сборки PTX, разбив 128-битные числа на 32-битные порции, вычислив 64-битные частичные произведения и добавив ихсоответственно.Очевидно, это займет немного работы.Можно получить разумную производительность на уровне C, разбив число на 64-битные куски и используя __umul64hi () в сочетании с обычным 64-битным умножением и некоторыми дополнениями.Это приведет к следующему:
__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;
}
Ниже приведена версия 128-разрядного умножения, использующего встроенную сборку PTX.Для этого требуется PTX 3.0, который поставляется с CUDA 4.2, а для кода требуется графический процессор с как минимум вычислительной способностью 2.0, то есть устройство класса Fermi или Kepler.В коде используется минимальное количество инструкций, поскольку для реализации 128-битного умножения необходимы шестнадцать 32-битных умножений.Для сравнения, приведенный выше вариант с использованием встроенных функций CUDA компилируется в 23 инструкции для цели 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;
}