सीयूडीए

Question

में 2 डी साझा मेमोरी की व्यवस्था कैसे की जाती है मैंने हमेशा रैखिक साझा मेमोरी (लोड, स्टोर, पड़ोसियों तक पहुंच) के साथ काम किया है, लेकिन मैंने बैंक संघर्षों का अध्ययन करने के लिए 2 डी में एक सरल परीक्षण किया है जिसके परिणामस्वरूप मुझे भ्रमित कर दिया गया है।

अगला कोड एक आयामी वैश्विक मेमोरी सरणी से साझा स्मृति में डेटा पढ़ता है और इसे साझा स्मृति से वैश्विक स्मृति में कॉपी करता है।

__global__ void update(int* gIn, int* gOut, int w) { 

// shared memory space 
__shared__ int shData[16][16]; 
// map from threadIdx/BlockIdx to data position 
int x = threadIdx.x + blockIdx.x * blockDim.x; 
int y = threadIdx.y + blockIdx.y * blockDim.y; 
// calculate the global id into the one dimensional array 
int gid = x + y * w; 

// load shared memory 
shData[threadIdx.x][threadIdx.y] = gIn[gid]; 
// synchronize threads not really needed but keep it for convenience 
__syncthreads(); 
// write data back to global memory 
gOut[gid] = shData[threadIdx.x][threadIdx.y]; 
} 

दृश्य प्रोफाइलर साझा स्मृति में संघर्ष की सूचना दी। अगले कोड से बचने के thouse संघर्ष (केवल अंतर दिखाने)

// load shared memory 
shData[threadIdx.y][threadIdx.x] = gIn[gid]; 

// write data back to global memory 
gOut[gid] = shData[threadIdx.y][threadIdx.x]; 

यह व्यवहार मुझे क्योंकि प्रोग्रामिंग व्यापक समानांतर प्रोसेसर में उलझन में है। एक हाथ से दृष्टिकोण हम पढ़ सकते हैं:

सी और सीयूडीए में मैट्रिक्स तत्व पंक्ति प्रमुख सम्मेलन के अनुसार रैखिक रूप से संबोधित स्थानों में रखा गया है। यही है, मैट्रिक्स की पंक्ति 0 के तत्वों को पहले लगातार स्थानों में क्रमशः रखा जाता है।

क्या यह साझा स्मृति व्यवस्था से संबंधित है? या धागे इंडेक्स के साथ? शायद मैं कुछ याद कर रहा हूँ?

गिरी विन्यास इस प्रकार है:

// kernel configuration 
dim3 dimBlock = dim3 (16, 16, 1); 
dim3 dimGrid = dim3 (64, 64); 
// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threads 
update<<<dimGrid, dimBlock>>>(d_input, d_output, 1024); 

अग्रिम धन्यवाद।

Answer 1

हां, साझा स्मृति को आपके द्वारा अपेक्षित पंक्ति-प्रमुख क्रम में व्यवस्थित किया गया है। तो अपने [16] [16] सरणी पंक्ति बुद्धिमान संग्रहीत किया जाता है, कुछ इस तरह:

 bank0 .... bank15 
row 0 [ 0 .... 15 ] 
    1 [ 16 .... 31 ] 
    2 [ 32 .... 47 ] 
    3 [ 48 .... 63 ] 
    4 [ 64 .... 79 ] 
    5 [ 80 .... 95 ] 
    6 [ 96 .... 111 ] 
    7 [ 112 .... 127 ] 
    8 [ 128 .... 143 ] 
    9 [ 144 .... 159 ] 
    10 [ 160 .... 175 ] 
    11 [ 176 .... 191 ] 
    12 [ 192 .... 207 ] 
    13 [ 208 .... 223 ] 
    14 [ 224 .... 239 ] 
    15 [ 240 .... 255 ] 
     col 0 .... col 15 

क्योंकि वहाँ 16 32 बिट साझा पूर्व फर्मी हार्डवेयर पर स्मृति बैंक हैं, प्रत्येक स्तंभ में हर पूर्णांक प्रविष्टि एक साझा पर नक्शे मेमोरी बैंक तो यह इंडेक्सिंग योजना की अपनी पसंद के साथ कैसे सहभागिता करता है?

ध्यान में रखना यह बात है कि ब्लॉक के भीतर धागे कॉलम प्रमुख क्रम के बराबर गिने जाते हैं (तकनीकी रूप से संरचना का एक्स आयाम सबसे तेज़ बदलता है, उसके बाद y, उसके बाद z)। तो जब आप इस अनुक्रमण योजना का उपयोग करें: एक आधा ताना भीतर

shData[threadIdx.x][threadIdx.y] 

धागे उसी स्तंभ, जो एक ही साझा स्मृति बैंक से पढ़ने का मतलब से पढ़ने हो जाएगा, और बैंक संघर्ष हो जाएगा। आप विपरीत योजना का उपयोग करते हैं: एक ही आधा ताना भीतर

shData[threadIdx.y][threadIdx.x] 

धागे एक ही पंक्ति है, जो 16 अलग साझा स्मृति बैंकों में से प्रत्येक से पढ़ने का मतलब से पढ़ कर दिया जाएगा, कोई संघर्ष होते हैं।