我刚刚在Linux Ubuntu 10.04下安装了我的cuda SDK。我的图形卡是NVIDIA GeForce GT 425M,我想用它来解决一些重大的计算问题。 我想知道的是:有没有办法使用一些无符号的128位int var?当使用gcc在CPU上运行我的程序时,我使用__uint128_t类型,但将它与cuda一起使用似乎不起作用。 有什么我可以做的在cuda上有128位整数?cuda上的128位整数?
非常感谢您 利玛窦蒙蒂 Msoft编程
我刚刚在Linux Ubuntu 10.04下安装了我的cuda SDK。我的图形卡是NVIDIA GeForce GT 425M,我想用它来解决一些重大的计算问题。 我想知道的是:有没有办法使用一些无符号的128位int var?当使用gcc在CPU上运行我的程序时,我使用__uint128_t类型,但将它与cuda一起使用似乎不起作用。 有什么我可以做的在cuda上有128位整数?cuda上的128位整数?
非常感谢您 利玛窦蒙蒂 Msoft编程
为了获得最佳性能,一个要映射在合适的CUDA矢量类型,诸如uint4的顶部的128位型,并且使用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位的部分乘积,并适当地将它们相加来构造乘法。显然这需要一些工作。有人可能会通过将数字分成64位块并使用__umul64hi()与常规的64位乘法和一些附加功能相结合来在C级获得合理的性能。这将导致以下结果:
__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;
}
下面是使用PTX内联汇编的128位乘法的一个版本。它需要随CUDA 4.2一起提供的PTX 3.0,并且代码要求至少具有计算能力2.0的GPU,即费米或开普勒类设备。该代码使用最少数量的指令,因为需要16个32位乘法来实现128位乘法。相比之下,上述使用CUDA内在函数的变体针对sm_20目标编译为23条指令。
__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;
}
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;
}
非常感谢!还有一个问题:从效率的角度来看,这是否足够快? – 2011-05-28 18:59:13
我测试了我的CPU上的代码。它实际上工作,但它比使用__uint128_t类型慢6倍...是否有任何方法使其更快? – 2011-05-28 22:04:45
你用CPU上的'my_uint128'在CPU上测试了内置的128位整数?当然,本地支持将会更快。我们希望这种128位类型的GPU的性能会比内置128位整数的CPU的性能更快。 – tkerwin 2011-05-28 22:52:46
非常感谢你!这正是我需要的! – 2011-11-03 14:12:34