atomicAdd() GPU

Question

मैं GPU पर एक परियोजना कर रहा हूँ पर डबल के लिए, और मैं डबल के लिए atomicAdd(), उपयोग करने के लिए क्योंकि CUDA डबल के लिए इसका समर्थन नहीं करता है, तो मैं नीचे दिए गए कोड है, जो NVIDIA प्रदान का उपयोग ।

__device__ double atomicAdd(double* address, double val) 
{ 
    unsigned long long int* address_as_ull = 
              (unsigned long long int*)address; 
    unsigned long long int old = *address_as_ull, assumed; 
    do { 
     assumed = old; 
     old = atomicCAS(address_as_ull, assumed, 
         __double_as_longlong(val + 
         __longlong_as_double(assumed))); 
    } while (assumed != old); 
    return __longlong_as_double(old); 
} 

अब मुझे पता है क्यों, एक पाश की आवश्यकता होती है, जबकि लागू करना चाहते हैं (मान लिया! = पुराने)

Answer 1

मूल रूप से कार्यान्वयन एक लोड है, जो atomically नहीं किया जा सकता आवश्यकता है क्योंकि। तुलना और स्वैप आपरेशन कोई गारंटी नहीं address पर मूल्य चक्र है कि मूल्य address से भरी हुई है और चक्र atomicCAS कॉल है कि बीच में परिवर्तन नहीं होगा नहीं है

(*address == assumed) ? (assumed + val) : *address 

की एक परमाणु संस्करण है अद्यतन मूल्य को स्टोर करने के लिए प्रयोग किया जाता है। यदि ऐसा होता है, तो address पर मान अपडेट नहीं किया जाएगा। इसलिए पाश सुनिश्चित करता है कि दो आपरेशन तक वहां पढ़ने और तुलना करें और स्वैप आपरेशन, जिसका मतलब है कि अद्यतन जगह ले ली के बीच address पर मूल्य का कोई बदलाव नहीं आया है दोहराया जाता है।