सीयूडीए ने वैश्विक मेमोरी तक पहुंच का सहारा लिया

Question

मैंने सीयूडीए प्रोग्रामिंग गाइड पढ़ा है, लेकिन मुझे एक चीज़ याद आई। आइए मान लें कि मेरे पास वैश्विक मेमोरी में 32 बिट इंटेल है और मैं इसे सहबद्ध एक्सेस के साथ साझा मेमोरी में कॉपी करना चाहता हूं। ग्लोबल सरणी में 0 से 1024 तक इंडेक्स हैं, और मान लें कि मेरे पास 256 धागे वाले प्रत्येक 4 ब्लॉक हैं।

__shared__ int sData[256]; 

कोलेसेड एक्सेस कब किया जाता है? वैश्विक स्मृति में

1.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y]; 

पतों ताना में 32 धागे से 255 0, प्रत्येक से नकल कर रहे हैं, इसलिए यहाँ यह ठीक है?

2.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex]; 

तो someIndex 32 के कई यह एकत्रित नहीं कर रहा है ना? Misaligned adresses? क्या वो सही है?

Answer 1

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

Answer 2

यदि आप 1 डी ग्रिड और थ्रेड-ज्यामिति का उपयोग करना चाहते हैं तो आपके उदाहरण सही हैं। मुझे लगता है कि आप जिस इंडेक्सिंग का उपयोग करना चाहते हैं वह [blockIdx.x*blockDim.x + threadIdx.x] है।

# 1 के साथ, 32 धागे एक वार्प में निष्पादित करते हैं कि निर्देश 'एक साथ' ताकि उनके अनुरोध, जो अनुक्रमिक हैं और 128 बी (32 x 4) के साथ गठबंधन हैं, मुझे विश्वास है कि टेस्ला और फर्मि वास्तुकला दोनों में सहवास है।

# 2 के साथ, यह थोड़ा धुंधला है। यदि someIndex 1 है, तो यह एक ताना में 32 अनुरोधों के सभी सम्मिलित नहीं होगा, लेकिन यह आंशिक कोलेसिंग कर सकता है। मेरा मानना ​​है कि फर्मि डिवाइसेस मेमोरी के 128 बी अनुक्रमिक खंड (और पहले 4 बी, जिसे थ्रेड की जरूरत नहीं है, बर्बाद कर दिया गया है) के हिस्से के रूप में एक तार में थ्रेड 1-31 के लिए पहुंच को जोड़ देगा। मुझे लगता है कि टेस्ला आर्किटेक्चर डिवाइसेस गलत तरीके से गलत तरीके से पहुंचने के कारण एक अनौपचारिक पहुंच बनाएगा, लेकिन मुझे यकीन नहीं है।

someIndex के साथ, 8, टेस्ला के पास 32 बी संरेखित पते होंगे, और फर्मि उन्हें 32 बी, 64 बी और 32 बी के रूप में समूहित कर सकता है। लेकिन लब्बोलुआब यह है, है someIndex और स्थापत्य कला के मूल्य पर निर्भर करता है, क्या होता है धुँधली है, और यह जरूरी भयानक नहीं होगा।

Answer 3

1 पर आपका अनुक्रमण गलत (या जानबूझकर इतनी अजीब यह गलत लगता है) है, कुछ ब्लॉक एक सूत्र में एक ही तत्व का उपयोग, इसलिए इन ब्लॉकों में एकत्रित पहुँच के लिए कोई रास्ता नहीं है।

सबूत:

उदाहरण:

Grid = dim(2,2,0) 

t(blockIdx.x, blockIdx.y) 

//complete block reads at 0 
t(0,0) -> sData[threadIdx.x] = gData[0]; 
//complete block reads at 2 
t(0,1) -> sData[threadIdx.x] = gData[2]; 
//definetly coalesced 
t(1,0) -> sData[threadIdx.x] = gData[threadIdx.x]; 
//not coalesced since 2 is no multiple of a half of the warp size = 16 
t(1,1) -> sData[threadIdx.x] = gData[threadIdx.x + 2]; 

तो इसकी एक "लक" खेल यदि एक ब्लॉक एकत्रित कर रहा है सामान्य में ऐसा है, तो कोई

लेकिन एकत्रित स्मृति पढ़ता नियम हैं पहले के रूप में नए cuda संस्करणों पर सख्त नहीं है।
लेकिन संगतता के मुद्दों के लिए आपको सबसे कम CUDA संस्करणों के लिए कर्नेल अनुकूलन करने के लिए, अगर यह संभव हो सकता है की कोशिश करनी चाहिए।

http://mc.stanford.edu/cgi-bin/images/0/0a/M02_4.pdf

Answer 4

क्या आप अंततः चाहते हैं कि क्या आपके ग्रिड और ब्लॉक -1 डी या 2 डी कर रहे हैं अपने इनपुट डेटा एक -1 डी या 2 डी सरणी है इस पर निर्भर करता है, और:

यहाँ कुछ अच्छा स्रोत है। सबसे आसान मामला 1 डी:

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x]; 

यह सहसंबंधित है। अंगूठे का नियम मैं उपयोग करता हूं कि ब्लॉक ऑफसेट (ब्लॉकडिम * ब्लॉकआईडीएक्स) के ऑफसेट के रूप में सबसे तेजी से भिन्न समन्वय (थ्रेडआईडीएक्स) जोड़ा जाता है। अंत परिणाम यह है कि ब्लॉक में धागे के बीच इंडेक्सिंग की तरफ 1 है। यदि घुमाव बड़ा हो जाता है, तो आप कोलेसिंग खो देते हैं।

सरल नियम (फर्मि और बाद में जीपीयू पर) यह है कि यदि एक थ्रू में सभी धागे के पते एक ही गठबंधन 128-बाइट रेंज में आते हैं, तो एक मेमोरी लेनदेन परिणाम देगा (माना जाता है कि भार के लिए कैशिंग सक्षम है , जो डिफ़ॉल्ट है)। यदि वे दो गठित 128-बाइट श्रेणियों में आते हैं, तो दो मेमोरी लेनदेन परिणाम इत्यादि।

GT2xx और पहले GPUs पर, यह अधिक जटिल हो जाता है। लेकिन आप प्रोग्रामिंग गाइड में इसका विवरण पा सकते हैं।

अतिरिक्त उदाहरण:

एकत्रित नहीं:

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x]; 

एकत्रित नहीं है, लेकिन GT200 पर और बाद में बहुत बुरा नहीं:

stride = 2; 
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x]; 

नहीं पाए सब पर:

stride = 32; 
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x]; 

कोआ lesced, 2 डी ग्रिड, 1 डी ब्लॉक:

int elementPitch = blockDim.x * gridDim.x; 
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
          blockIdx.x * blockDim.x + threadIdx.x]; 

एकत्रित, 2 डी ग्रिड और ब्लॉक:

int x = blockIdx.x * blockDim.x + threadIdx.x; 
int y = blockIdx.y * blockDim.y + threadIdx.y; 
int elementPitch = blockDim.x * gridDim.x; 
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];