👨‍👩‍👦 👨🏾‍🤝‍👨🏻 🧗🏿 NVIDIA CUDA में साझा बैंक संघर्षों को समझना 🚑 🤽🏾 🚠

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

साझा मेमोरी संघर्ष कैसे उत्पन्न होते हैं?

टकराव तब होता है जब एक ताना (डिवाइस संस्करण 2.0 के लिए) से 2 या अधिक धाराएं या ताना का आधा (डिवाइस संस्करण 1.3 और निम्न के लिए) एक्सेस बाइट्स जो एक ही मेमोरी बैंक में स्थित विभिन्न 32 बिट शब्दों से संबंधित हैं। संघर्ष के मामले में, पहुंच अनुक्रमिक है। बैंक तक पहुँचने वाले थ्रेड्स की संख्या को संघर्ष की डिग्री कहा जाता है। यदि संघर्ष की डिग्री एन है, तो पहुंच एन से धीमी है यदि कोई संघर्ष नहीं था।

प्रसारण तंत्र

संस्करण 1.x के उपकरणों पर, एक संघर्ष से बचा जा सकता है यदि कई धागे एक ही बैंक से संबंधित एक ही शब्द का उपयोग करते हैं, और केवल अगर यह अनुरोध एकल है - इस मामले में, प्रसारण एक्सेस तंत्र का उपयोग किया जाता है।

संस्करण 2.x के उपकरणों पर, ऐसे कई अनुरोध हो सकते हैं और उन्हें समानांतर में लागू किया जाएगा (विभिन्न थ्रेड्स शब्द के विभिन्न बाइट्स तक पहुंच सकते हैं)।

उपकरणों के संस्करण 2.0 पर पहुंच सुविधाएँ

64 बिट एक्सेस के साथ, एक बैंक संघर्ष केवल तभी होता है, जब किसी भी ताना के आधे हिस्से में से 2 या उससे अधिक धाराएँ एक ही बैंक से संबंधित पते तक पहुँचती हैं।

128-बिट एक्सेस के साथ, आमतौर पर दूसरी डिग्री के बैंकों का टकराव होता है।

32 से अधिक क्षमता के साथ प्रवेश 32, 64 और 128 बिट्स की क्षमता के साथ अनुरोधों में टूट गया है।

बैंकों को कैसे मेमोरी आवंटित की जाती है

मेमोरी को बैंकों के बीच इस तरह वितरित किया जाता है कि अनुक्रम में प्रत्येक 32 बिट शब्द क्रमिक रूप से उपकरण संस्करण 2.0 के मामले में 32 बैंकों में से एक और डिवाइस संस्करण 1.3 और निम्न के मामले में 16 बैंकों को सौंपा जाता है। तदनुसार, बैंक संख्या की गणना निम्न सूत्र द्वारा की जा सकती है:

बैंक नंबर = (बाइट्स में पता / 4)% 32 - डिवाइस संस्करण 2.0 के लिए

बैंक नंबर = (बाइट्स में पता / 4)% 16 - डिवाइस संस्करण 1.x के लिए

संघर्ष एक्सेस उदाहरण

उपकरणों के लिए संस्करण 1.x

1.8 और 16 बिट एक्सेस

__shared__ char shmem8[32]; 
      

        
        
        
      

     char data = shmem8[threadIdx.x];

इस उदाहरण में, पहले 4 बाइट्स एक ही बैंक में हैं, इसलिए पहले 4 बाइट्स एक्सेस पर विरोध करेंगे

समस्या को अनावश्यक डेटा (पैडिंग) जोड़कर और एक्सेस योजना को बदलकर हल किया जाता है:

__shared__ char shmem8[32*4]; 
      

        
        
        
      

     char data = shmem8[threadIdx.x*4];

16-बिट एक्सेस के लिए:

__shared__ short shmem16[32]; 
      

        
        
        
      

     short data = shmem16[threadIdx.x];

इस उदाहरण में, पहले 2 शॉर्ट्स एक ही बैंक में हैं, इसलिए पहले 2 प्रवाह एक्सेस पर संघर्ष करेंगे

समस्या को 8-बिट एक्सेस के समान हल किया गया है:

__shared__ short shmem16[32*2]; 
      

        
        
        
      

     short data = shmem16[threadIdx.x*2];

2. 32-बिट एक्सेस

इस प्रकार की पहुंच के लिए, बैंक संघर्ष कम स्पष्ट हैं, लेकिन उदाहरण के लिए, यह एक्सेस स्कीम हो सकती है:

__shared__ int shmem32[64]; 
      

        
        
        
      

     int data1 = shmem32[threadIdx.x*2]; 
      

        
        
        
      

     int data2 = shmem32[threadIdx.x*2+1];

इस मामले में, 0 और 8 वीं धारा क्रमशः 0 और 1 बैंकों से पढ़ी जाती है, इस प्रकार 2 डिग्री का संघर्ष पैदा होता है।

उदाहरण के लिए, आप इस समस्या को इस तरह हल कर सकते हैं:

__shared__ int shmem32_1[32]; 
      

        
        
        
      

     __shared__ int shmem32_2[32]; 
      

        
        
        
      

     int data1 = shmem32_1[threadIdx.x]; 
      

        
        
        
      

     int data2 = shmem32_2[threadIdx.x];

उपकरणों के लिए संस्करण 2.0

प्रसारण एक्सेस की सुविधाओं के कारण, इन उपकरणों पर 8 और 16 बिट एक्सेस योजनाएं बैंक टकराव का कारण नहीं बनती हैं, हालांकि, निम्नलिखित मामले में संघर्ष हो सकता है:

__shared__ int shared[64]; 
      

        
        
        
      

     int data = shared[threadIdx.x*s];

यदि विरोध समाप्‍त होता है तो भी। यदि एस विषम है, लेकिन कोई संघर्ष नहीं होता है।

बैंक संघर्ष ट्रैकिंग

NVIDIA बैनक परीक्षक

CUT_BANK_CHECKER (सरणी, अनुक्रमणिका) मैक्रो का उपयोग करके विरोधों को ट्रैक किया जा सकता है, जो CUDA सुविधा टूलकिट का हिस्सा है। ऐसा करने के लिए, आपको इस मैक्रो का उपयोग मेमोरी तक पहुंचने और इम्यूलेशन मोड में एप्लिकेशन को चलाने के लिए करना होगा। जब आवेदन पूरा हो जाता है, तो एक संघर्ष रिपोर्ट मुद्रित की जाएगी।

उदाहरण के लिए, इस तरह:

__shared__ int shared[64]; 
      

        
        
        
      

     int data = CUT_BANK_CHECKER(shared, threadIdx.x*s);

CUDA प्रोफाइलर

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

निष्कर्ष

अंत में, मैं ध्यान देता हूं कि बैंक संघर्षों को हल करने के लिए सबसे प्रभावी तकनीक एक्सेस योजनाओं को विकसित करना है जो कि उनकी घटना को कम कर देता है और प्रोफाइलर द्वारा आवेदन के बाद के विश्लेषण (जो कभी भी अतिशयोक्तिपूर्ण नहीं है)।

NVIDIA CUDA में साझा बैंक संघर्षों को समझना