128-битное целое число на cuda?

Мне просто удалось установить мой cuda SDK под Linux Ubuntu 10.04. Моя графическая карта – NVIDIA GeForce GT 425M, и я бы хотел использовать ее для некоторых сложных вычислительных задач. Я задаюсь вопросом: есть ли способ использовать некоторые unsigned 128 bit int var? При использовании gcc для запуска моей программы на CPU я использовал тип __uint128_t, но использование его с cuda, похоже, не работает. Есть ли что-нибудь, что я могу сделать, чтобы иметь 128 бит целых чисел на cuda?

Большое спасибо Matteo Monti Msoft Programming

Для лучшей производительности хотелось бы сопоставить 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-битного умножения, в которой используется встроенная assembly PTX. Он требует PTX 3.0, который поставляется вместе с CUDA 4.2, а для этого кода требуется графический процессор с минимальной вычислительной способностью 2.0, то есть устройством classа Fermi или Kepler. В коде используется минимальное количество инструкций, так как для реализации 128-битного умножения требуется шестнадцать 32-битных умножений. Для сравнения, вариант, используемый с использованием CUDA intrinsics, составляет 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"(ax), "r"(ay), "r"(az), "r"(aw), "r"(bx), "r"(by), "r"(bz), "r"(bw)); return res; } 

CUDA не поддерживает 128-битные целые числа. Вы можете подделать операции самостоятельно, используя два 64-битных целых числа.

Посмотрите на это сообщение :

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

Очень запоздалый ответ, но вы могли бы рассмотреть возможность использования этой библиотеки:

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

который определяет структуру размером 128 бит, с методами и автономными функциями полезности, чтобы заставить ее функционировать должным образом, что позволяет использовать ее как регулярное целое число. В основном.

Давайте будем гением компьютера.