क्यूडा पर 128 बिट पूर्णांक?

Question

मैं बस लिनक्स उबंटू 10.04 के तहत अपना क्यूडा एसडीके स्थापित करने में कामयाब रहा। मेरा ग्राफिक कार्ड एक एनवीआईडीआईए जेफफोर्स जीटी 425 एम है, और मैं इसे कुछ भारी कम्प्यूटेशनल समस्या के लिए उपयोग करना चाहता हूं। मुझे आश्चर्य है कि: क्या कुछ हस्ताक्षरित 128 बिट int var का उपयोग करने का कोई तरीका है? सीपीयू पर अपना प्रोग्राम चलाने के लिए जीसीसी का उपयोग करते समय, मैं __uint128_t प्रकार का उपयोग कर रहा था, लेकिन इसे cuda के साथ उपयोग करना प्रतीत नहीं होता है। क्या कोई ऐसा कुछ है जो मैं क्यूडा पर 128 बिट पूर्णांक रखने के लिए कर सकता हूं?

आपको बहुत बहुत Matteo मोंटि Msoft प्रोग्रामिंग

Answer 1

सर्वश्रेष्ठ प्रदर्शन के लिए, ऐसे ही एक uint4 के रूप में एक उपयुक्त CUDA वेक्टर प्रकार, के शीर्ष पर 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; 
} 

गुणा इसी तरह, 32-बिट टुकड़ों में 128 बिट संख्या को तोड़ने की गणना 64-बिट आंशिक उत्पादों और उन्हें उचित रूप से जोड़कर PTX इनलाइन विधानसभा का प्रयोग कर बनाया जा सकता है। जाहिर है यह थोड़ा सा काम लेता है। 64-बिट भाग में संख्या को तोड़कर और नियमित 64-बिट गुणा और कुछ जोड़ों के साथ संयोजन में __umul64hi() का उपयोग करके सी स्तर पर उचित प्रदर्शन हो सकता है। यह निम्नलिखित में परिणाम होगा:

__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 बिट गुणा का एक संस्करण है। इसके लिए पीटीएक्स 3.0 की आवश्यकता है, जो सीयूडीए 4.2 के साथ भेज दिया गया है, और कोड को कम से कम गणना क्षमता 2.0, यानी फर्मि या केप्लर क्लास डिवाइस के साथ एक जीपीयू की आवश्यकता होती है। कोड कम से कम निर्देशों का उपयोग करता है, क्योंकि 128-बिट गुणा को लागू करने के लिए सोलह 32-बिट गुणों की आवश्यकता होती है। तुलनात्मक रूप से, 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; 
}