सीयूडीए डॉट उत्पाद

Question

मैं विभिन्न ब्लॉकों में अंतिम योग की परमाणु गणना के साथ डबल परिशुद्धता सरणी के लिए क्लासिक डॉट-उत्पाद कर्नेल को कार्यान्वित करने का प्रयास कर रहा हूं। मैंने प्रोग्रामिंग गाइड के पेज 116 में बताए गए डबल परिशुद्धता के लिए परमाणु का उपयोग किया। संभवतः मैं कुछ गलत कर रहा हूं। प्रत्येक ब्लॉक में धागे के आंशिक रकम सही ढंग से गणना की जाती हैं लेकिन बाद में परमाणु ऑपरेशन ठीक से काम नहीं कर रहा है चूंकि हर बार जब मैं एक ही डेटा के साथ अपना कर्नेल चलाता हूं, तो मुझे अलग-अलग परिणाम मिलते हैं। अगर कोई गलती तलाश सकता है या वैकल्पिक समाधान प्रदान कर सकता है तो मैं आभारी रहूंगा! यहाँ

__global__ void cuda_dot_kernel(int *n,double *a, double *b, double *dot_res) 
{ 
    __shared__ double cache[threadsPerBlock]; //thread shared memory 
    int global_tid=threadIdx.x + blockIdx.x * blockDim.x; 
    int i=0,cacheIndex=0; 
    double temp = 0; 
    cacheIndex = threadIdx.x; 
    while (global_tid < (*n)) { 
     temp += a[global_tid] * b[global_tid]; 
     global_tid += blockDim.x * gridDim.x; 
    } 
    cache[cacheIndex] = temp; 
    __syncthreads(); 
    for (i=blockDim.x/2; i>0; i>>=1) { 
     if (threadIdx.x < i) { 
      cache[threadIdx.x] += cache[threadIdx.x + i]; 
     } 
     __syncthreads(); 
    } 
    __syncthreads(); 
    if (cacheIndex==0) { 
     *dot_res=cuda_atomicAdd(dot_res,cache[0]); 
    } 
} 

और मेरे डिवाइस समारोह atomicAdd है: यहाँ मेरी गिरी है

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

    return old; 
} 

Answer 1

आप cuda_atomicAdd फ़ंक्शन गलत तरीके से उपयोग कर रहे हैं। आपके कर्नेल का यह अनुभाग:

if (cacheIndex==0) { 
    *dot_res=cuda_atomicAdd(dot_res,cache[0]); 
} 

अपराधी है। यहां, आप परमाणु रूप से dot_res में जोड़ें। तो गैर परमाणुdot_res सेट के परिणामस्वरूप यह लौटाता है। इस फ़ंक्शन से रिटर्न परिणाम परमाणु रूप से अपडेट किए जाने वाले स्थान के पिछले मान है, और यह केवल "जानकारी" या कॉलर के स्थानीय उपयोग के लिए प्रदान किया गया है। आप इसे असाइनिक रूप से अपडेट किए गए कार्यों को असाइन नहीं करते हैं, जो पहले स्थान पर परमाणु स्मृति पहुंच का उपयोग करने के उद्देश्य को पूरी तरह से हरा देता है।इसके बजाए ऐसा कुछ करें:

if (cacheIndex==0) { 
    double result=cuda_atomicAdd(dot_res,cache[0]); 
} 

Answer 2

कमी तदर्थ CUDA कोड का उपयोग कर मुश्किल हो सकता है सही हो रही है, तो यहां कोई वैकल्पिक समाधान एक थ्रस्ट कलन विधि का उपयोग है , जिसे CUDA टूलकिट के साथ शामिल किया गया है:

#include <thrust/inner_product.h> 
#include <thrust/device_ptr.h> 

double do_dot_product(int *n, double *a, double *b) 
{ 
    // wrap raw pointers to device memory with device_ptr 
    thrust::device_ptr<double> d_a(a), d_b(b); 

    // inner_product implements a mathematical dot product 
    return thrust::inner_product(d_a, d_a + n, d_b, 0.0); 
} 

Answer 3

गहराई से आपके कोड की जांच नहीं की गई लेकिन यहां कुछ सलाह दी गई हैं।
मैं केवल जोर देने के लिए सलाह दूंगा यदि आप केवल अपने जीपीयू का उपयोग ऐसे सामान्य कार्यों के लिए करते हैं, क्योंकि यदि कोई जटिल समस्या उत्पन्न होगी तो लोगों को gpu पर समानांतर प्रोग्राम करने का कोई विचार नहीं है।

1. डॉट उत्पाद को सारांशित करने के लिए एक नया समांतर कमी कर्नेल प्रारंभ करें।
   चूंकि डेटा पहले से ही डिवाइस पर है, इसलिए आपको एक नया कर्नेल शुरू करने में प्रदर्शन में कमी दिखाई नहीं देगी।

2. आपका कर्नेल नवीनतम GPU पर संभावित ब्लॉक की अधिकतम संख्या में स्केल नहीं करना प्रतीत होता है। यदि ऐसा होता है और आपका कर्नेल लाखों मूल्यों के डॉट उत्पाद की गणना करने में सक्षम होगा तो धारावाहिक परमाणु संचालन के कारण प्रदर्शन नाटकीय रूप से घट जाएगा।

3. शुरुआती गलती: क्या आपका इनपुट डेटा और साझा मेमोरी एक्सेस रेंज चेक किया गया है? या आप सुनिश्चित हैं कि इनपुट डेटा हमेशा आपके ब्लॉक आकार का एकाधिक होता है? अन्यथा आप कचरा पढ़ेंगे। मेरे अधिकांश गलत परिणाम इस गलती के कारण थे।

4. अपनी समांतर कमी को अनुकूलित करें। My Thesis या Optimisations Mark Harris

untested है, मैं सिर्फ यह नीचे लिखा नोटपैड में:

/* 
* @param inCount_s unsigned long long int Length of both input arrays 
* @param inValues1_g double* First value array 
* @param inValues2_g double* Second value array 
* @param outDots_g double* Output dots of each block, length equals the number of blocks 
*/ 
__global__ void dotProduct(const unsigned long long int inCount_s, 
    const double* inValuesA_g, 
    const double* inValuesB_g, 
    double* outDots_g) 
{ 
    //get unique block index in a possible 3D Grid 
    const unsigned long long int blockId = blockIdx.x //1D 
      + blockIdx.y * gridDim.x //2D 
      + gridDim.x * gridDim.y * blockIdx.z; //3D 


    //block dimension uses only x-coordinate 
    const unsigned long long int tId = blockId * blockDim.x + threadIdx.x; 

    /* 
    * shared value pair products array, where BLOCK_SIZE power of 2 
    * 
    * To improve performance increase its size by multiple of BLOCK_SIZE, so that each threads loads more then 1 element! 
    * (outDots_g length decreases by same factor, and you need to range check and initialize memory) 
    * -> see harris gpu optimisations/parallel reduction slides for more informations. 
    */ 
    __shared__ double dots_s[BLOCK_SIZE]; 


    /* 
    * initialize shared memory array and calculate dot product of two values, 
    * shared memory always needs to be initialized, its never 0 by default, else garbage is read later! 
    */ 
    if(tId < inCount_s) 
     dots_s[threadIdx.x] = inValuesA_g[tId] * inValuesB_g[tId]; 
    else 
     dots_s[threadIdx.x] = 0; 
    __syncthreads(); 

    //do parallel reduction on shared memory array to sum up values 
    reductionAdd(dots_s, dots_s[0]) //see my thesis link 

    //output value 
    if(threadIdx.x == 0) 
     outDots_g[0] = dots_s[0]; 

    //start new parallel reduction kernel to sum up outDots_g! 
} 

संपादित करें: हटाया अनावश्यक अंक।