cuda上的128位整数?

我只是设法在Linux Ubuntu 10.04下安装我的cuda SDK。 我的图形卡是NVIDIA GeForce GT 425M,我想用它来解决一些重大的计算问题。 我想知道的是:有没有办法使用一些无符号的128位int var? 当使用gcc在CPU上运行我的程序时,我使用__uint128_t类型,但将它与cuda一起使用似乎不起作用。 有什么我可以做的在cuda上有128位整数?

非常感谢Matteo Monti Msoft编程


为获得最佳性能,人们希望将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;nt"
         "addc.cc.u32     %1, %5, %9;nt"
         "addc.cc.u32     %2, %6, %10;nt"
         "addc.u32        %3, %7, %11;nt"
         : "=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 ("{nt"
         "mul.lo.u32      %0, %4, %8;    nt"
         "mul.hi.u32      %1, %4, %8;    nt"
         "mad.lo.cc.u32   %1, %4, %9, %1;nt"
         "madc.hi.u32     %2, %4, %9,  0;nt"
         "mad.lo.cc.u32   %1, %5, %8, %1;nt"
         "madc.hi.cc.u32  %2, %5, %8, %2;nt"
         "madc.hi.u32     %3, %4,%10,  0;nt"
         "mad.lo.cc.u32   %2, %4,%10, %2;nt"
         "madc.hi.u32     %3, %5, %9, %3;nt"
         "mad.lo.cc.u32   %2, %5, %9, %2;nt"
         "madc.hi.u32     %3, %6, %8, %3;nt"
         "mad.lo.cc.u32   %2, %6, %8, %2;nt"
         "madc.lo.u32     %3, %4,%11, %3;nt"
         "mad.lo.u32      %3, %5,%10, %3;nt"
         "mad.lo.u32      %3, %6, %9, %3;nt"
         "mad.lo.u32      %3, %7, %8, %3;nt"
         "}"
         : "=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;
} 
链接地址: http://www.djcxy.com/p/45767.html

上一篇: 128 bit integer on cuda?

下一篇: What does GCC