Cuda'da 128 bit tam sayı?

Question

Linux Ubuntu 10.04'ün altında cuda SDK'yı yüklemeyi başardım. Grafik kartım bir NVIDIA geForce GT 425M'dir ve bazı ağır hesaplama problemleri için kullanmak istiyorum. Şaşırtıcı olan: bazı imzasız 128 bit int var kullanmak için herhangi bir yolu var mı? Programımı CPU üzerinde çalıştırmak için gcc kullanırken, __uint128_t türünü kullanıyordum, ancak cuda ile çalışmak işe yaramıyor. Cuda'da 128 bit tam sayıya sahip olmak için yapabileceğim bir şey var mı?

size doğal 128 bitlik tamsayılar desteklemez

Answer 1

, böyle bir uint4 gibi uygun CUDA vektör tipi, üst üste 128-bit türü harita, ve PTX satır içi derleme kullanarak işlevselliğini uygulamak isteyeyim . ek şöyle olacaktır:

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

benzer şekilde 32 bit parçalar halinde 128 bit numaraları kırma 64 bit kısmi ürünlerin işlem ve uygun ekleyerek PTX içi düzeneği kullanılarak inşa edilebilir çarpma. Açıkçası bu biraz iş gerektiriyor. C seviyesinde 64 bitlik parçalara bölünerek ve normal 64-bit çarpma ve bazı eklemelerle uyumlu olarak __umul64hi() kullanılarak C seviyesinde makul bir performans elde edilebilir. Bu aşağıda neden olacaktır:

Aşağıda

__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 satır içi montaj kullanır 128 bit çarpma bir sürümüdür. CUDA 4.2 ile birlikte gönderilen PTX 3.0'ı gerektirir ve bu kod, en azından hesaplama kapasitesi 2.0 olan bir GPU, yani bir Fermi veya Kepler sınıfı cihaz gerektirir. Kod, 128 bitlik bir çoğaltmayı uygulamak için on altı adet 32 ​​bitlik çarpma gerektiğinden, en az sayıda komut kullanır. Karşılaştırma yapmak gerekirse, yukarıda CUDA intrinsics kullanan varyant, sm_20 hedefi için 23 talimatı derler.

__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'yı Programlama çok Matteo Monti Msoft ederiz. İki 64 bit tam sayı kullanarak işlemleri kendiniz yapabilirsiniz.

this post de

Görünüş:

En iyi performans için

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