cuda上的128位整数?
Posted
技术标签:
【中文标题】cuda上的128位整数?【英文标题】:128 bit integer on cuda? 【发布时间】:2011-09-03 23:08:42 【问题描述】:我刚刚设法在 Linux Ubuntu 10.04 下安装了我的 cuda SDK。我的显卡是 NVIDIA geForce GT 425M,我想用它来解决一些繁重的计算问题。 我想知道的是:有没有办法使用一些无符号的 128 位 int var?当使用 gcc 在 CPU 上运行我的程序时,我使用的是 __uint128_t 类型,但是将它与 cuda 一起使用似乎不起作用。 有什么办法可以在 cuda 上使用 128 位整数吗?
【问题讨论】:
【参考方案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;
通过将 128 位数字分解为 32 位块、计算 64 位部分乘积并将它们适当地相加,可以类似地使用 PTX 内联汇编来构造乘法。显然,这需要一些工作。通过将数字分解为 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,即 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;
【讨论】:
@njuffa - 我想今天你会建议一个基于 2 个 64 位值的解决方案? @einpoklum 不太可能,因为 64 位整数运算是模拟的,通常最好在本机指令而不是其他模拟之上构建模拟。因为 32 位整数乘法和乘法加法本身在 Maxwell 和 Pascal 架构上进行仿真,所以最好使用本地 16 位 乘法,它映射到机器指令XMAD
(a 16x16+32 位乘加运算)。我读到,原生 32 位整数乘法是使用 Volta 架构恢复的,但我还没有使用 Volta 的实践经验。
与 32 位整数相比性能如何? 1/16 还是类似的?
@huseyintugrulbuyukisik 根据指令计数,它大约是原生 32 位乘法的 1/16。根据功能单元的加载和寄存器使用情况,实际的性能影响可能会有所不同,具体取决于代码上下文。
我们也可以原子地添加 uint128 吗?【参考方案2】:
CUDA 本身不支持 128 位整数。您可以使用两个 64 位整数自己伪造操作。
看this 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;
【讨论】:
非常感谢!还有一个问题:从效率的角度来看,这是否足够快? 我在我的 CPU 上测试了该代码。它确实有效,但它比使用 __uint128_t 类型慢 6 倍......没有什么办法让它更快吗? 您使用 CPU 上的my_uint128
测试了 CPU 上的内置 128 位整数?当然原生支持会更快。希望这种 128 位类型的 GPU 上的性能将比内置 128 位整数的 CPU 上的性能更快。
链接坏了吗?【参考方案3】:
一个迟来的答案,但你可以考虑使用这个库:
https://github.com/curtisseizert/CUDA-uint128
它定义了一个 128 位大小的结构,具有方法和独立的实用程序函数以使其按预期运行,从而使其可以像常规整数一样使用。大部分。
【讨论】:
这真的很酷,而且比其他答案要好得多:) 在查看源代码后,我看到有一个 __mul64hi PTX 指令可以使 64 * 64 位乘法高效。以上是关于cuda上的128位整数?的主要内容,如果未能解决你的问题,请参考以下文章
如何从 16 x 8 位 __m128i 值中提取 32 x 4 位整数