128 bit integer na cuda?
Właśnie udało mi się zainstalować mój cuda SDK pod Linuksem Ubuntu 10.04. Moja karta graficzna to NVIDIA geForce GT 425M i chciałbym jej użyć w przypadku jakiegoś ciężkiego problemu obliczeniowego. Zastanawiam się, czy jest jakiś sposób, aby użyć jakiegoś unsigned 128 bit int var? Podczas używania gcc do uruchamiania mojego programu na procesorze, używałem typu _ _ uint128 _ T, ale używanie go z cuda nie wydaje się działać. Czy jest coś, co mogę zrobić, aby mieć 128 bitowe liczby całkowite na cuda?
Dziękuję bardzo Matteo Monti Msoft Programowanie
3 answers
Aby uzyskać najlepszą wydajność, warto odwzorować 128-bitowy typ na odpowiednim typie wektorowym CUDA, takim jak uint4, i zaimplementować funkcjonalność za pomocą wbudowanego montażu PTX. Dodatek wyglądałby mniej więcej tak:
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;
}
Mnożenie może być podobnie skonstruowane za pomocą wbudowanego montażu PTX, dzieląc 128-bitowe liczby na 32-bitowe kawałki, obliczając 64-bitowe produkty częściowe i dodając je odpowiednio. Oczywiście wymaga to trochę pracy. Można dostać rozsądna wydajność na poziomie C poprzez podzielenie liczby na 64-bitowe kawałki i użycie _ _ umul64hi () w połączeniu ze zwykłym 64-bitowym mnożeniem i pewnymi dodatkami. Wynikałoby to z:
__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;
}
Poniżej znajduje się Wersja 128-bitowego mnożenia wykorzystującego wbudowany montaż PTX. Wymaga PTX 3.0, który jest dostarczany z CUDA 4.2, A kod wymaga GPU z co najmniej zdolnością obliczeniową 2.0, czyli urządzenia klasy Fermi lub Kepler. Kod wykorzystuje minimalną liczbę instrukcje, ponieważ do implementacji mnożenia 128-bitowego potrzebne jest szesnaście 32-bitowych mnożenia. Dla porównania, powyższy wariant przy użyciu CUDA intrinsics kompiluje się do 23 instrukcji dla celu 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;
}
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/doraprojects.net/template/agent.layouts/content.php on line 54
2012-07-21 03:00:41
CUDA nie obsługuje natywnie 128-bitowych liczb całkowitych. Możesz sfałszować operacje używając dwóch 64-bitowych liczb całkowitych.
Spójrz na ten 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;
}
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/doraprojects.net/template/agent.layouts/content.php on line 54
2011-05-28 15:28:44
Odpowiedź o wiele spóźniona, ale mógłbyś rozważyć użycie tej biblioteki:
Https://github.com/curtisseizert/CUDA-uint128
, która definiuje strukturę o wielkości 128 bitów, z metodami i wolnostojącymi funkcjami użytkowymi, aby uzyskać jej działanie zgodnie z oczekiwaniami, co pozwala jej używać jak zwykłej liczby całkowitej. Głównie.
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/doraprojects.net/template/agent.layouts/content.php on line 54
2018-05-30 13:00:17