cuda上的128位整数？

Question

我刚刚成功在 Linux Ubuntu 10.04 下安装了我的 cuda SDK。我的显卡是 NVIDIA geForce GT 425M，我想用它来解决一些繁重的计算问题。 我想知道的是：有没有办法使用一些无符号的 128 位 int var？当使用 gcc 在 CPU 上运行我的程序时，我使用的是 __uint128_t 类型，但将它与 cuda 一起使用似乎不起作用。 有什么办法可以让 cuda 上有 128 位整数吗？

Answer 1

为了获得最佳性能，人们希望将 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 位部分乘积并适当地相加。显然这需要一些工作。通过将数字分解为 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 位乘法的版本。它需要 PTX 3.0（随 CUDA 4.2 一起提供），并且代码需要至少具有计算能力 2.0 的 GPU，即 Fermi 或 Kepler 类设备。该代码使用最少数量的指令，因为需要 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;
}

Answer 2

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

Answer 3

对于后代，请注意，从 11.5 开始，当主机编译器支持（例如，clang/gcc，但不支持 MSVC）时，CUDA 和 nvcc 在设备代码中支持 __int128_t。 11.6 添加了对 __int128_t 调试工具的支持。

请参阅：

• https://developer.nvidia.com/ blog/cuda-11-6-toolkit-new-release-revealed/
• https://developer.nvidia.com/blog/implementing-high- precision-decimal-arithmetic-with-cuda-int128/

Answer 4

一个迟来的答案，但您可以考虑使用这个库：

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

定义了一个128位大小的结构，具有方法和独立的实用函数来使其按预期运行，这使得它可以像常规整数一样使用。大多。