2012-02-08 17 views
15

मैं इस प्रकार आव्यूह गुणन के लिए एक चर घोषित करने के लिए कोशिश कर रहा हूँ:Cuda साझा मेमोरी सरणी चर

__shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 

मैं यह इतना उपयोगकर्ता इनपुट मैट्रिक्स के आकार की गणना करने के लिए कर सकता है बनाने की कोशिश कर रहा हूँ कि होगा हालांकि, मतलब BLOCK_SIZE बदल रहा है। मैंने इसे बदल दिया लेकिन मुझे एक कंपाइलर त्रुटि मिल रही है: "त्रुटि: निरंतर मान ज्ञात नहीं है"। मैंने इसे देखा है और यह thread के समान है। इसलिए मैं करने की कोशिश की:

__shared__ int buf []; 

लेकिन तब मैं मिलता है: "त्रुटि: अधूरा प्रकार अनुमत नहीं है"

धन्यवाद, कोड के साथ दान अद्यतन (काफी this guide पीछा किया और CUDA गाइड के साथ बाहर घूर) : ब्लॉक आकार को मैट्रिक्स के आकार के उपयोगकर्ता से पूछकर पारित किया जाता है। वे एक्स और वाई दर्ज करते हैं। ब्लॉक आकार केवल एक्स है और अभी इसे एक्स और वाई के समान आकार को स्वीकार करना है।

__global__ void matrixMul(float* C, float* A, float* B, int wA, int wB,size_t block_size) 
{ 
    // Block index 
    int bx = blockIdx.x; 
    int by = blockIdx.y; 

    // Thread index 
    int tx = threadIdx.x; 
    int ty = threadIdx.y; 

    // Index of the first sub-matrix of A processed 
    // by the block 
    int aBegin = wA * block_size * by; 

    // Index of the last sub-matrix of A processed 
    // by the block 
    int aEnd = aBegin + wA - 1; 

    // Step size used to iterate through the 
    // sub-matrices of A 
    int aStep = block_size; 

    // Index of the first sub-matrix of B processed 
    // by the block 
    int bBegin = block_size * bx; 

    // Step size used to iterate through the 
    // sub-matrices of B 
    int bStep = block_size * wB; 
    float Csub=0; 
    // Loop over all the sub-matrices of A and B 
    // required to compute the block sub-matrix 
    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    { 
     // Declaration of the shared memory array As 
     // used to store the sub-matrix of A 

     extern __shared__ float As[]; 

     // Declaration of the shared memory array Bs 
     // used to store the sub-matrix of B 
     extern __shared__ float Bs[]; 
     extern __shared__ float smem[]; 

     // Load the matrices from global memory 
     // to shared memory; each thread loads 
     // one element of each matrix 
     smem[ty*block_size+tx] = A[a + wA * ty + tx]; 
     //cuPrintf("\n\nWhat are the memory locations?\n"); 
     //cuPrintf("The shared memory(A) is: %.2f\n",smem[ty*block_size+tx]); 
     smem[block_size*block_size+ty*block_size+tx] = B[b + wB * ty + tx]; 
     //cuPrintf("The shared memory(B) is: %.2f\n",smem[block_size*block_size+ty*block_size+tx]); 
     // Synchronize to make sure the matrices 
     // are loaded 
     __syncthreads(); 

     // Multiply the two matrices together; 
     // each thread computes one element 
     // of the block sub-matrix 
     for (int k = 0; k < block_size; ++k) 
     { 

      Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ; 
      //cuPrintf("Csub is currently: %.2f\n",Csub); 
     } 
     //cuPrintf("\n\n\n"); 
     // Synchronize to make sure that the preceding 
     // computation is done before loading two new 
     // sub-matrices of A and B in the next iteration 
     //cuPrintf("the results are csub: %.2f\n",Csub); 
     __syncthreads(); 
    } 
    // Write the block sub-matrix to device memory; 
    // each thread writes one element 
    int c = wB * block_size * by + block_size * bx; 
    C[c + wB * ty + tx] = Csub; 


} 
+0

धागा का उपयोग आप एक गलती निहित से जुड़ा हुआ। – talonmies

+0

आह मैं देखता हूं कि आपने इसे संपादित किया, धन्यवाद। – Dan

उत्तर

26

extern __shared__ int buf[];

जब आप कर्नेल का शुभारंभ आप इसे इस तरह से लॉन्च किया जाना चाहिए;

kernel<<<blocks,threads,numbytes_for_shared>>>(...);

आप साझा की कई extern घोषणा है:

extern __shared__ float As[];

extern __shared__ float Bs[];

इस AsBs रूप में एक ही पते की ओर इशारा करते के लिए नेतृत्व करेंगे।

आपको 1 डी-सरणी के अंदर As और Bs रखना होगा।

extern __shared__ float smem[]; 

जब गिरी बुला, आप इसे 2*BLOCK_SIZE*BLOCK_SIZE*sizeof(float) साथ प्रारंभ होना चाहिए।

जब रूप में का अनुक्रमण, smem[y*BLOCK_SIZE+x] का उपयोग करें और जब बी एस में का अनुक्रमण smem[BLOCK_SIZE*BLOCK_SIZE+y*BLOCK_SIZE+x]

+0

धन्यवाद जो काम करता है, लेकिन मैं अब अन्य त्रुटियों में चल रहा हूं – Dan

+0

जानकारी के लिए धन्यवाद। मुझे अभी मैट्रिक्स को गुणा करने के साथ थोड़ा सा मुद्दा है, लेकिन यहां दी गई जानकारी के साथ इसका प्रयास करेंगे। – Dan

+0

मैं इसे लूप में कैसे संभाल सकता हूं जहां ए और बी को के साथ एक्सेस किया जा रहा है? मेरे पास यही है: (int k = 0; k Dan

0

सही लगता है।

आम तौर पर इस मामले में आपको कुछ मॉलोक करने की आवश्यकता होगी।

यहां दो चीजें हैं, एक सी को 2 डी एरे के बारे में पता नहीं है (यह केवल सरणी की एक सरणी है) और सरणी आकारों को समय स्थिरांक (या संकलक समय पर संकलन कर सकते हैं) को संकलित करने की आवश्यकता होती है।

यदि आप सी 99 का उपयोग कर रहे हैं तो आप फ़ंक्शन के पैरामीटर का उपयोग करके सरणी आकार घोषित कर सकते हैं, लेकिन सी 99 समर्थन ... सबसे अच्छा है।

+0

मैंने mallocing की कोशिश की है, लेकिन मुझे विश्वास नहीं है कि आपको डिवाइस कोड पर ऐसा करने की अनुमति है। – Dan

+0

डिवाइस कोड पर 'malloc' को कॉल नहीं कर सकता ... कर्नेल में प्रवेश करने से पहले सभी गतिशील मेमोरी आवंटित की जानी चाहिए, और गतिशील बफर को आवंटित करने की आवश्यकता है और 'malloc' के CUDA- विशिष्ट संस्करणों का उपयोग करके डिवाइस पर प्रतिलिपि बनाई जानी चाहिए और 'memcpy'। – Jason

+0

@ जेसन: वास्तव में, फर्मि जीपीयू पर, दोनों 'मॉलोक' और सी ++ 'नया' ऑपरेटर दोनों समर्थित हैं। लेकिन केवल आवंटन के लिए जो वैश्विक स्मृति में रहेंगे। आप यह कहते हुए सही हैं कि गतिशील रूप से आवंटित * ​​साझा * मेमोरी को कॉलिंग होस्ट कोड द्वारा आवंटित किया जाना चाहिए (इस मामले में कर्नेल लॉन्च सिंटैक्स के हिस्से के रूप में या एक अलग API कॉल के माध्यम से)। – talonmies

23

आपके पास कर्नेल - स्थैतिक या गतिशील के अंदर साझा स्मृति को घोषित करने के लिए दो विकल्प हैं। मुझे लगता है कि क्या आप इस समय क्या कर रहे हैं इस तरह दिखता है:

#define BLOCK_SIZE (16) 

__global__ void sgemm0(const float *A, const float *B, float *C) 
{ 
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 

} 

और आप आसानी से BLOCK_SIZE को बदलने में सक्षम होना चाहते हैं।

template<int blocksize=16> 
__global__ void sgemm1(const float *A, const float *B, float *C) 
{ 
    __shared__ float As[blocksize][blocksize]; 

} 
template void sgemm1<16>(const float *, const float *, float *C); 

तो फिर तुम आप के रूप में संकलन समय पर के रूप में कई अलग अलग ब्लॉक आकार वेरिएंट का दृष्टांत कर सकते हैं:

एक संभावना यह इस तरह स्थिर साझा स्मृति आवंटन उपयोग करने के लिए जारी है, लेकिन आवंटन आकार टेम्पलेट पैरामीटर बनाने के लिए, है जरुरत।

आप गतिशील, स्मृति को आबंटित इस तरह यह निर्धारित करना चाहते हैं:

__global__ void sgemm2(const float *A, const float *B, float *C) 
{ 
    extern __shared__ float As[]; 

} 

और फिर गिरी कॉल करने के लिए एक तर्क के रूप आवंटन के आकार जोड़ें:

size_t blocksize = BLOCK_SIZE * BLOCK_SIZE; 
sgemm2<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....); 

आप तो कई स्थैतिक रूप से घोषित एरे हैं जिन्हें आप गतिशील रूप से आवंटित साझा स्मृति के साथ प्रतिस्थापित करना चाहते हैं, फिर ध्यान रखें कि प्रति कर्नेल केवल एक गतिशील साझा स्मृति आवंटन है, इसलिए कई आइटम उस स्मृति खंड (शेयर) के भीतर निकलते हैं।तो अगर आप की तरह कुछ था:

#define BLOCK_SIZE (16) 

__global__ void sgemm0(const float *A, const float *B, float *C) 
{ 
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 

} 

आप से बदलने सकता है:

#define BLOCK_SIZE (16) 

__global__ void sgemm3(const float *A, const float *B, float *C) 
{ 
    extern __shared__ float buffer[]; 

    float *As = &buffer[0]; 
    float *Bs = &buffer[BLOCK_SIZE*BLOCK_SIZE]; 

} 

और इस तरह गिरी का शुभारंभ:

size_t blocksize = 2 * BLOCK_SIZE * BLOCK_SIZE; 
sgemm3<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....); 

सभी समान रूप से मान्य हैं, हालांकि मैं व्यक्तिगत रूप से पक्ष टेम्पलेट संस्करण क्योंकि यह स्वचालित लूप अनलोलिंग जैसे अन्य कंपाइलर अनुकूलन की अनुमति दे सकता है कि गतिशील संस्करण अतिरिक्त काम के बिना नहीं हो सकता है।

+0

'@talonmies, कि __shared__ फ्लोट * के रूप में; 'बाहरी __shared__ फ्लोट जैसा होना चाहिए [];' जैसा कि ब्रानो के उत्तर में है। आप दोनों को एक उथल-पुथल मिलती है। :) – harrism

+0

बह, उसको याद किया। धन्यवाद मार्क। – talonmies

+0

ठीक है उत्तर के लिए धन्यवाद। मैंने बाहरी साझा किया। हालांकि यह इसे 1 डी सरणी में परिवर्तित करता है, कार्यक्रम 2 डी सरणी का उपयोग कर मूल रूप से था। मैं समझता हूं कि 2 डी सरणी सरणी के सरणी हैं इसलिए 1 डी सरणी काम करनी चाहिए।उदाहरण के लिए: [ty] [tx] = ए [ए + डब्ल्यूए * टाई + टीएक्स] के रूप में; मैं -> [ty * MAX_THREADS + tx] = ए [ए + डब्ल्यूए * टाई + टीएक्स] में परिवर्तित किया गया; जहां maxthreads 1023 है, क्योंकि मेरे अधिकतम धागे 1024 हैं। लेकिन मुझे 2x2 मैट्रिक्स के लिए -0,0, -0,0 मिल रहा है। – Dan

संबंधित मुद्दे