cuda上的128位整數？

Question

我剛剛在Linux Ubuntu 10.04下安裝了我的cuda SDK。我的圖形卡是NVIDIA GeForce GT 425M，我想用它來解決一些重大的計算問題。 我想知道的是：有沒有辦法使用一些無符號的128位int var？當使用gcc在CPU上運行我的程序時，我使用__uint128_t類型，但將它與cuda一起使用似乎不起作用。 有什麼我可以做的在cuda上有128位整數？

非常感謝您 利瑪竇蒙蒂 Msoft編程

Answer 1

爲了獲得最佳性能，一個要映射在合適的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; 
} 

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