समानांतर कंप्यूटिंग टूल का उपयोग करते समय, यह बहुत संभावना है कि एक ऐसी स्थिति उत्पन्न हो सकती है जब एल्गोरिथ्म में दो ऐसे अनुक्रमिक चरण होते हैं: i ) प्रत्येक j- th स्ट्रीम कुछ मध्यवर्ती गणना परिणाम को j- th मेमोरी सेल में संग्रहीत करता है, और फिर, ii ) इस धारा को एक के परिणामों का उपयोग करना चाहिए या अधिक "पड़ोसी" धागे। स्पष्ट रूप से, प्रोग्राम कोड में एक निश्चित समय अवरोध को व्यवस्थित करना आवश्यक है, जो प्रत्येक थ्रेड को तब तक खत्म कर देता है जब सभी ने अपने मध्यवर्ती परिणामों को संबंधित मेमोरी कोशिकाओं (चरण ( i )) में सहेज लिया हो। अन्यथा, कुछ थ्रेड चरण ( ii ) पर जा सकते हैं, जबकि कुछ अन्य थ्रेड्स ने अभी तक चरण ( i ) को पूरा नहीं किया है। यह खेदजनक है, लेकिन CUDA के रचनाकारों ने महसूस किया कि किसी भी GPU पर किसी भी संख्या में थ्रेड को सिंक्रनाइज़ करने के लिए इस तरह के एक विशेष अंतर्निहित तंत्र की आवश्यकता नहीं है। तो आप इस संकट से कैसे निपट सकते हैं? हालाँकि, Google, प्रॉम्प्ट्स को देखते हुए, इस मुद्दे से परिचित है, यह अपने कार्य के लिए एक तैयार संतोषजनक नुस्खा खोजने के लिए संभव नहीं था, और शुरुआती (जो मैं हूं) के लिए वांछित परिणाम प्राप्त करने के रास्ते में कुछ नुकसान हैं।

CUDA वास्तुकला के बारे में कुछ शब्द

शुरू करने के लिए, मैं आपको आधिकारिक दस्तावेज [1,2] और स्लाइड्स [3,4] , विभिन्न तृतीय-पक्ष साइटों की सामग्री [5-11], सामान्य तस्वीर जो कि एक प्रोग्रामर का सामना करता है, CUDA का उपयोग करते समय आपको याद दिलाता हूं। अमूर्तता के उच्चतम स्तर पर, उसे SIMT ( सिंगल-इंस्ट्रक्शन, मल्टीपल-थ्रेड ) आर्किटेक्चर के साथ एक समानांतर कंप्यूटिंग सिस्टम मिलता है - एक कमांड को कई या कम स्वतंत्र थ्रेड के साथ समानांतर में निष्पादित किया जाता है। एकल कार्य (चित्र 1 देखें) के भाग के रूप में चलने वाले इन सभी धागों की समग्रता को ग्रिड कहा जाता है ।

अंजीर। 1।

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

अंजीर। 2।

उदाहरण के लिए, मेरे टेस्ला में 30 एसएम हैं, जिनमें से प्रत्येक में 8 स्केलर प्रोसेसर हैं। हालाँकि, इन 240 कोर पर, आप उपलब्ध संसाधनों को साझा करने के लिए हार्डवेयर तंत्रों की बदौलत बड़ी संख्या में थ्रेड्स ⁽¹⁾ से ग्रिड चला सकते हैं (इन कोर और उपलब्ध मेमोरी दोनों का कार्य समय)। और बस इन तंत्रों के कार्यान्वयन की कुछ विशेषताएं साझा स्मृति तक पहुँचने पर थ्रेड्स को सिंक्रनाइज़ करने की तकनीक निर्धारित करती हैं।

इस तरह की महत्वपूर्ण विशेषताओं में से एक ताना-बाना में 32 टुकड़ों के प्रवाह का समूह है, जो बड़े संरचनाओं - ब्लॉकों के हिस्से बनते हैं । प्रत्येक ब्लॉक के सभी धागे (उदाहरण के लिए, मेरे टेस्ला ब्लॉक में अधिकतम 512 थ्रेड ^{(1) हो सकते हैं} ) एक एसएम पर सख्ती से चलते हैं, इसलिए उनके पास केवल इसके संसाधनों तक पहुंच है। हालांकि, एक एसएम (एक छवि 3 देखें) पर एक से अधिक ब्लॉक लॉन्च किए जा सकते हैं, और संसाधनों को उनके बीच समान रूप से विभाजित किया जाएगा।

अंजीर। 3।

प्रत्येक एसएम में, एक नियंत्रण इकाई होती है जो प्रोसेसर समय संसाधन वितरित करती है। ऐसा इसलिए किया जाता है ताकि समय पर हर क्षण, एक एसएम के सभी कर्नेल एक ताना कड़ाई से निष्पादित करें। और इसके पूरा होने पर, इस एसएम को सौंपा गया अगला ताना चालाक तरीके से चुना जाता है। इस प्रकार, यह पता चला है कि CUDA की हार्डवेयर विशेषता के कारण एक ताना का प्रवाह सिंक्रनाइज़ है और इसे SIMD ( सिंगल इंस्ट्रक्शन, मल्टीपल डेटा ) के करीब एक विधि के अनुसार निष्पादित किया जाता है। लेकिन अलग-अलग ताना से एक ब्लॉक का प्रवाह भी सिंक से दूर हो सकता है।

एक और, कोई कम महत्वपूर्ण विशेषता CUDA में मेमोरी का संगठन और इसके विभिन्न हिस्सों तक धागे की पहुंच नहीं है। स्ट्रीम के लिए सामान्य पहुंच की उच्चतम डिग्री वैश्विक मेमोरी (भौतिक मेमोरी ) द्वारा प्रदान की जाती है, भौतिक रूप से ग्राफिक्स एडेप्टर बोर्ड पर सील किए गए एकीकृत सर्किट के रूप में कार्यान्वित की जाती है - वही वीडियो मेमोरी, जिसे अब गीगाबाइट में गणना की जाती है। प्रोसेसर के बाहर का स्थान इस प्रकार की मेमोरी को वीडियो कार्ड पर कंप्यूटिंग के लिए उपलब्ध अन्य की तुलना में सबसे धीमा बनाता है। एक छोटी "सुगमता की डिग्री" साझा की गई मेमोरी है : प्रत्येक एसएम (छवि 2 देखें) में स्थित ब्लॉक, आमतौर पर आकार में 16KB ⁽¹⁾ , केवल उन थ्रेड्स तक पहुंचता है जो इस एसएम के कोर पर चलते हैं (देखें अंजीर। 1, अंजीर। 3)। चूंकि एक एसएम पर समानांतर निष्पादन के लिए एक से अधिक ब्लॉक आवंटित किए जा सकते हैं, एसएम में उपलब्ध साझा मेमोरी की पूरी मात्रा इन ब्लॉकों के बीच समान रूप से वितरित की जाती है। यह ध्यान दिया जाना चाहिए कि साझा मेमोरी शारीरिक रूप से एसएम कोर के बहुत करीब कहीं स्थित है, इसलिए इसमें रजिस्टरों की गति, स्मृति के मुख्य प्रकार की तुलना में एक उच्च पहुंच गति है। यह रजिस्टर है जो प्राथमिक मशीन निर्देशों के ऑपरेंड के रूप में काम कर सकता है, और सबसे तेज़ मेमोरी है। इस एसएम पर चलने वाले सभी थ्रेड्स के बीच एक एसएम के सभी कैश रजिस्टर समान रूप से विभाजित होते हैं। किसी भी धागे द्वारा उपयोग के लिए आबंटित रजिस्टरों का एक समूह उसे और केवल उसे उपलब्ध है। CUDA की शक्ति के चित्रण के रूप में (या, इसके विपरीत, आपदा के पैमाने पर): एक ही टेस्ला में, प्रत्येक एसएम 32-बिट सामान्य-उद्देश्य रजिस्टर ^{(1) के} 16384 टुकड़ों का उपयोग करने के लिए प्रदान करता है।

उपरोक्त सभी से, हम यह निष्कर्ष निकाल सकते हैं कि एक ब्लॉक के प्रवाह के बीच की बातचीत को उनकी आम तेजी से साझा की गई मेमोरी के माध्यम से और दो अलग-अलग ब्लॉकों के प्रवाह के बीच - केवल वैश्विक मेमोरी का उपयोग करने की कोशिश की जानी चाहिए। यह वह जगह है जहाँ परिचय में उल्लिखित समस्या उत्पन्न होती है: स्मृति क्षेत्रों को पढ़ने और लिखने के लिए सार्वजनिक रूप से उपलब्ध विभिन्न धाराओं में डेटा की प्रासंगिकता पर नज़र रखना। दूसरे शब्दों में, थ्रेड सिंक्रोनाइज़ेशन की समस्या। जैसा कि पहले ही उल्लेख किया गया है, एक ब्लॉक के भीतर, प्रत्येक ताना के प्रवाह को एक दूसरे के साथ सिंक्रनाइज़ किया जाता है। ताना सदस्यता की परवाह किए बिना, ब्लॉक प्रवाह को सिंक्रनाइज़ करने के लिए, कई बाधा-प्रकार कमांड हैं:

__syncthreads () सबसे सुरक्षित तरीका है। यह फ़ंक्शन प्रत्येक थ्रेड को तब तक प्रतीक्षा करेगा जब तक (ए) इस ब्लॉक के अन्य सभी धागे इस बिंदु तक नहीं पहुंचते और (बी) इस ब्लॉक के थ्रेड्स द्वारा साझा किए गए और वैश्विक मेमोरी तक पहुंचने के सभी ऑपरेशन पूरे हो जाते हैं और इस ब्लॉक के थ्रेड्स के लिए दृश्यमान हो जाते हैं। यदि यह कथन है तो इस कमांड को सशर्त के अंदर रखना आवश्यक नहीं है, लेकिन ब्लॉक के सभी थ्रेड्स द्वारा इस फ़ंक्शन को बिना शर्त कॉल प्रदान करना आवश्यक है।
__threadfence_block () उस थ्रेड को बना देगा जो इसे प्रतीक्षा करता है जब तक कि साझा और वैश्विक मेमोरी तक सभी प्रतिबद्ध पहुंच संचालन पूरा नहीं हो जाता है और इस ब्लॉक के थ्रेड्स के लिए दृश्यमान हो जाते हैं।
__threadfence () उस थ्रेड को मजबूर करेगा जिसके कारण वह प्रतीक्षा करता है जब तक कि साझा मेमोरी के सभी प्रतिबद्ध एक्सेस ऑपरेशंस इस ब्लॉक के थ्रेड्स के लिए दृश्यमान न हो जाएं, और "डिवाइस" पर सभी थ्रेड्स के लिए वैश्विक मेमोरी वाले ऑपरेशंस। "डिवाइस" से मतलब ग्राफिक्स एडॉप्टर से है।
__threadfence_system () __threadfence () के समान है, लेकिन यह सीपीयू ("होस्ट") पर थ्रेड्स के साथ सिंक्रनाइज़ेशन को सक्षम करता है जब बहुत सुविधाजनक पेज-लॉक मेमोरी का उपयोग किया जाता है। [१,२] में और अधिक विवरण और कुछ अन्य स्रोत नीचे दी गई सूची में सूचीबद्ध हैं।

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

अधिक कार्य करें

इसलिए, अधिक विशेष रूप से, निम्नलिखित उदाहरण पर विचार करें। एडॉप्टर की वैश्विक मेमोरी में दो खंड आवंटित किए जाएं: X [] और P [] सरणियों के लिए 128 तत्व। सरणी X [] होस्ट से लिखा गया है (कंप्यूटर के रैम से केंद्रीय प्रोसेसर द्वारा)। 64 थ्रेड्स के दो ब्लॉकों का एक ग्रिड बनाएं - प्रत्येक में, कुल 128 थ्रेड्स (चित्र 4 देखें)।

अंजीर। 4।

अब चरण ( i ) का प्रदर्शन किया जा सकता है: नंबर j के साथ प्रत्येक स्ट्रीम, सरणी [[] के सभी तत्वों को एक-दूसरे से जोड़ देगा, जिसके परिणामस्वरूप P [j] लिखा जाएगा। अगला, चरण ( ii ) निष्पादित किया जाना चाहिए: प्रत्येक j- th स्ट्रीम, सरणी P [] के सभी तत्वों का योग शुरू करेगी, उन्हें संबंधित X [j] में लिख देगी। बेशक, समानांतर में 128 बार एक ही चीज़ को निष्पादित करने के लिए CUDA का उपयोग करना व्यर्थ है, लेकिन वास्तविक जीवन में प्रत्येक थ्रेड में भारित कारकों का अपना सेट होगा जिसके साथ योग होता है, और रूपांतरण एक्स -> पी और इसके विपरीत, पी > एक्स - कई बार होगा। हमारे उदाहरण में, हम एकता के बराबर गुणांक चुनते हैं - स्पष्टता और सादगी के लिए, जो सामान्यता का उल्लंघन नहीं करता है।

हम सिद्धांत से प्रयोग करने के लिए गुजरते हैं। एल्गोरिथ्म बहुत पारदर्शी है, और एक व्यक्ति जिसने कभी मल्टीथ्रेडिंग से निपटा नहीं है वह तुरंत निम्नलिखित CUDA कर्नेल कोड का सुझाव दे सकता है:

__global__ void Kernel(float *X, float *P) { const int N = 128; //       . const int index = threadIdx.x + blockIdx.x*blockDim.x; //  . float a; //   .   . /*  (i): */ a = X[0]; for(int j = 1; j < N; ++j) // ,   a += X[j]; P[index] = a / N; // ,     . /*   (i). */ /*  (ii): */ a = P[0]; for(int j = 1; j < N; ++j) // ,   a += P[j]; X[index] = a / N; // ,     . /*   (ii). */ }

इस कर्नेल के बार-बार निष्पादन से पता चलता है कि सरणी P [] समय-समय पर समान होगा, लेकिन, यहाँ, X [] कभी-कभी अलग हो सकता है। इसके अलावा, अगर कोई अंतर है, तो यह एक तत्व X [j] में नहीं होगा, बल्कि 32 लगातार तत्वों के समूह में होगा! इस मामले में, त्रुटिपूर्ण ब्लॉक में पहले तत्व का सूचकांक भी 32 का एक गुणक होगा - यह केवल उन बहुत ताना-बाना में से कुछ में सिंक्रनाइज़ेशन का एक प्रकटीकरण है और विभिन्न ताना`ओव की सिंक धाराओं में से एक है। यदि किसी थ्रेड में कोई त्रुटि हुई, तो यह उसके बाकी के सभी ताना में होगा। यदि आप CUDA डेवलपर्स द्वारा प्रस्तावित सिंक्रोनाइज़ेशन तंत्र को लागू करते हैं

 __global__ void Kernel(float *X, float *P) { ... /*   (i). */ __syncthreads(); /*  (ii): */ ... }

तब हम प्राप्त करेंगे कि प्रत्येक ब्लॉक स्ट्रीम समान परिणाम देगा। और अगर कहीं यह गलत होगा - तो पूरे ब्लॉक। इस प्रकार, यह किसी भी तरह से अलग-अलग ब्लॉक को सिंक्रनाइज़ करने के लिए रहता है।

समाधान के तरीके

दुर्भाग्य से, मैं केवल दो तरीकों को जानता हूं:

CUDA कर्नेल समाप्त हो जाता है और यदि केवल सभी थ्रेड समाप्त हो जाते हैं। इस प्रकार, एक कोर को दो में विभाजित किया जा सकता है और मुख्य कार्यक्रम से क्रमिक रूप से बुलाया जा सकता है;
वैश्विक स्मृति में झंडे की एक प्रणाली के साथ आओ।

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

एक ध्वज प्रणाली के साथ दूसरे विकल्प के रूप में, एक समान तंत्र का उल्लेख "बी .5 मेमोरी फ़ेंस फ़ंक्शंस" [1] अनुभाग में किया गया है। हालांकि, थोड़ा अलग कर्नेल एल्गोरिथ्म वहां माना जाता है। ब्लॉक सिंक्रोनाइज़ेशन को लागू करने के लिए, हम दो कार्यों को शुरू करते हैं: पहला खर्च किए गए ब्लॉकों के काउंटर के मूल्यों को तैयार करेगा, और दूसरा एक अवरोध की भूमिका निभाएगा - प्रत्येक थ्रेड को विलंबित करें जब तक कि सभी ब्लॉक पूरा न हो जाएं। उदाहरण के लिए, ये फ़ंक्शन और कर्नेल जो उनका उपयोग करते हैं, वे इस तरह दिख सकते हैं:

 __device__ unsigned int count; // -  .    //4     . /*    -: */ __device__ void InitSyncWholeDevice(const int index) { if (index == 0) //    grid` ( 0)   count = 0; //    . if (threadIdx.x == 0) //    block`  ,  - while (count != 0); //   . //     block` ,      : __syncthreads(); // , - .    device -  . } /*     device: */ __device__ void SyncWholeDevice() { //      : unsigned int oldc; //   ,     gmem  smem,    grid`: __threadfence(); //    block`    (  ) //-: if (threadIdx.x == 0) { //  oldc   count  "+1": oldc = atomicInc(&count, gridDim.x-1); //   ,    ""    gmem: __threadfence(); //     (   count      ), //     count,    ,    //  gmem.    ,    "",      //,      ""   . if (oldc != (gridDim.x-1)) while (count != 0); } //      ,      : __syncthreads(); } __global__ void Kernel_Synced(float *X, float *P) { InitSyncWholeDevice(threadIdx.x + blockIdx.x*blockDim.x); ... /*   (i). */ SyncWholeDevice(); /*  (ii): */ ... }

वह, जैसे, वह सब। झंडे को जख्मी कर दिया गया था, कार्य लिखे गए थे। यह पहले और दूसरे तरीकों के प्रदर्शन की तुलना करने के लिए बनी हुई है। लेकिन, दुर्भाग्यवश, SyncWholeDevice () फ़ंक्शन काउंटर को बढ़ाएगा , लेकिन यह एक बाधा देरी प्रदान नहीं करेगा। यह प्रतीत होता है, क्यों? सब के बाद, एक समय लूप है । यहाँ हम अमूर्त में उल्लिखित गद्यांश की ओर नौकायन कर रहे हैं, जो दिखाई देता है: यदि आप संकलक [12-14] द्वारा उत्पन्न nvcc ptx फाइल को देखते हैं , तो यह पता चलता है कि वह अपने दृष्टिकोण से एक खाली पाश फेंकता है। आप कंपाइलर को कम से कम दो तरीकों से लूप को इस तरह से ऑप्टिमाइज़ न करने के लिए मजबूर कर सकते हैं।

Ptx असेंबलर में एक स्पष्ट प्रविष्टि निश्चित रूप से काम करेगी। उदाहरण के लिए, ऐसा फ़ंक्शन, जिसके कॉल को लूप को बदलना चाहिए:

 __device__ void do_while_count_not_eq(int val) { asm("{\n\t" "$my_while_label: \n\t" " .reg .u32 r_count; \n\t" " .reg .pred p; \n\t" " ld.global.u32 r_count, [count]; \n\t" " setp.ne.u32 p, r_count, %0; \n\t" "@p bra $my_while_label; \n\t" "}\n\t" : : "r"(val)); }

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

 __device__ volatile unsigned int count; // -  .    //4     . ... //  oldc   count  "+1": oldc = atomicInc((unsigned int*)&count, gridDim.x-1); ...

समाधान विधियों का आकलन

अब हम दो ब्लॉक तुल्यकालन विधियों के प्रदर्शन का एक मोटा सैद्धांतिक अनुमान लगाते हैं। अफवाह यह है कि एक कर्नेल कॉल ~ 10 माइक्रोसेकंड लेता है - यह कई कोर कॉल द्वारा सिंक्रनाइज़ेशन की कीमत है। लूप से एक बाधा को पेश करके सिंक्रनाइज़ेशन के मामले में, ~ 10 थ्रेड (कितने ब्लॉक पर निर्भर करता है) वेतन वृद्धि और लूप में वैश्विक मेमोरी में एक सेल पढ़ता है, जहां प्रत्येक इनपुट / आउटपुट ऑपरेशन में लगभग 500 घड़ी चक्र लगते हैं। प्रत्येक ब्लॉक को इस तरह के संचालन को करने दें 3. फिर, लगभग 10 * 500 * 3 = 1.5 * 10 ^ 4 चक्रों को सिंक्रनाइज़ेशन ऑपरेशन पर खर्च किया जाएगा। 1.5 GHz की मुख्य आवृत्ति पर, हमें 1.0 * 10 ^ (- 5) सेकंड = 10 μs मिलते हैं। अर्थात परिमाण का क्रम समान है।

लेकिन, ज़ाहिर है, कम से कम कुछ परीक्षणों के परिणामों को देखना दिलचस्प है। अंजीर। 5 में, पोस्ट रीडर 100 लगातार रूपांतरण एक्स -> पी -> एक्स पर बिताए समय की तुलना देख सकता है, प्रत्येक ग्रिड कॉन्फ़िगरेशन के लिए 10 बार दोहराया गया। 100 परिवर्तनों ^{(2) के} लिए आवश्यक समय को औसत करने के लिए 10 बार दोहराना होता है।

अंजीर। 5।

क्षैतिज तल में, शुरुआती ब्लॉक की संख्या और उनमें से प्रत्येक में थ्रेड्स की संख्या प्लॉट की जाती है। " मल्टी कर्नेल लॉन्च " विधि ( MKL ) के सापेक्ष " सिंगल कर्नेल लॉन्च " विधि के लिए समय अक्ष को " अक्षीय कॉल ( एसकेएल - सिंगल कॉलनेल लॉन्च) कहते हैं । यह स्पष्ट रूप से देखा जाता है कि विचाराधीन ग्रिड कॉन्फ़िगरेशन पर लाभ, हालांकि बहुत छोटा है, लगभग हमेशा सकारात्मक होता है। हालांकि, जितने अधिक ब्लॉक हैं, उतना कम एमकेएल विधि प्रदर्शन में पिछड़ जाती है। 32 ब्लॉक के लिए, वह SKL विधि से थोड़ा बाहर निकलता है। यह इस तथ्य के कारण है कि अधिक ब्लॉक, अधिक थ्रेड्स ( थ्रेडआईडीएक्सएक्सएक्सएक्स == 0 होने पर) धीमी वैश्विक मेमोरी से गिनती चर पढ़ते हैं। लेकिन कोई तंत्र नहीं है "मैंने इसे एक बार पढ़ा, सभी प्रवाह को अर्थ दिया"। यदि हम एक ब्लॉक में थ्रेड्स की संख्या के आधार पर स्वयं की निरंतर संख्या के आधार पर सापेक्ष उत्पादकता में बदलाव पर विचार करते हैं, तो आप कुछ नियमितता को भी नोटिस कर सकते हैं। लेकिन यहां, लेखक के काम के लिए अज्ञात प्रभाव, ब्लॉक में प्रवाह के सिंक्रनाइज़ेशन से संबंधित है, एसएम में ताना प्रबंधन। इसलिए, हम आगे की टिप्पणियों से बचते हैं।

यह काम करने वाले धागे (1024) के समान संख्या के साथ प्रदर्शन को देखना दिलचस्प है, लेकिन ब्लॉकों में उनका अलग विभाजन। चित्रा 6 दो तरीकों (एमकेएल और एसकेएल) के लिए उपरोक्त परिवर्तनों के 100 * 10 पर बिताए गए असामान्य समय के रेखांकन को दर्शाता है।

अंजीर। 6।

वास्तव में, यह चित्र 5 में विकर्ण "टुकड़ा" है। यह स्पष्ट रूप से देखा जाता है कि सबसे पहले, बड़े ब्लॉकों के साथ, दोनों सिंक्रनाइज़ेशन विधियों का प्रदर्शन समान रूप से बढ़ता है। CUDA डेवलपर्स ने आधिकारिक दस्तावेज [2] में इस तरह के प्रभाव की चेतावनी दी है, लेकिन लेखक, फिर से, दुर्भाग्य से, इस घटना के तंत्र का विवरण नहीं जानता है। अंतर की कमी और यहां तक कि ब्लॉक में सबसे छोटे विभाजन के साथ एसकेएल विधि का नुकसान जुड़ा हुआ है, जैसा कि पहले ही उल्लेख किया गया है, चर गिनती की रीडिंग की संख्या में वृद्धि के साथ।

यह ध्यान दिया जाना चाहिए कि परीक्षण पीकेएक्स-कोडांतरक सम्मिलित के साथ लूप को प्रतिस्थापित करके SKL विधि के कार्यान्वयन के दौरान किए गए थे। कभी-कभी अस्थिर स्पेसियर का उपयोग करना (ग्रिड कॉन्फ़िगरेशन के आधार पर) प्रक्रिया को धीमा कर देता है, और कभी-कभी इसे गति देता है। मंदी का परिमाण 0.20% तक पहुंच जाता है, और त्वरण 0.15% है। यह व्यवहार, सबसे अधिक संभावना है, संकलक द्वारा और एक व्यक्ति द्वारा ptx- कोडांतरक सम्मिलित करते समय लूप के कार्यान्वयन की सुविधाओं से निर्धारित होता है, और हमें SKL विधि के दोनों कार्यान्वयन को समान रूप से उत्पादक बनाने पर विचार करने की अनुमति देता है।

निष्कर्ष

इस लेख में, मैंने एक बुनियादी स्तर पर थ्रेड सिंक्रोनाइज़ेशन, ब्लॉक सिंक्रोनाइज़ेशन के तरीकों की समस्या पर रोशनी डालने की कोशिश की; कुछ परीक्षणों के बाद, चित्रों को CUDA प्रणाली का एक सामान्य विवरण दें। इसके अलावा, परीक्षण कार्यक्रम ⁽²⁾ के स्रोत कोड में ^, पाठक साझा मेमोरी में बफर के विश्वसनीय उपयोग का एक और उदाहरण पा सकेंगे (थ्रेड्स __syncthreads () ) के माध्यम से सिंक्रनाइज़ किए जाते हैं। मुझे आशा है कि किसी को यह उपयोगी लगेगा। व्यक्तिगत रूप से, यह जानकारी, एक जगह एकत्र की गई, कोड के साथ प्रयोग करने के कई दिनों को बचाने में मदद करेगी, और "googling", क्योंकि मुझे दस्तावेज़ीकरण पढ़ने के लिए बहुत चौकस नहीं होने की एक बेवकूफ प्रवृत्ति है।

⁽¹⁾ यह कंप्यूटर पर उपलब्ध एडेप्टर के बारे में तकनीकी जानकारी प्राप्त करने के लिए CUDA API फ़ंक्शन cudaGetDeviceProperties (...) [1-2] का उपयोग करने का प्रस्ताव है।
⁽²⁾ परीक्षण कार्यक्रम के लिए स्रोत कोड , pastebin.com पर अपलोड किया गया।

सूचना के स्रोतों की सूची

[१] क्यूडा सी प्रोग्रामिंग गाइड
[२] क्यूडा सी बेस्ट प्रैक्टिस गाइड
[३] उन्नत CUDA वेबिनार: मेमोरी ऑप्टिमाइज़ेशन
[४] एस। तारिक, जीपीयू कम्प्यूटिंग और सीयूडीएए आर्किटेक्चर का एक परिचय
[५] वेंडरबिल्ट यूनिवर्सिटी, ACCRE, CUDA के साथ GPU कम्प्यूटिंग
[६] ओएमएसटीयू, रेडियो इंजीनियरिंग संकाय, एकीकृत सूचना संरक्षण विभाग, जीपीआर के लिए प्रोग्राम " रिटर्मिंग "
[[] ग्रीष्मकालीन सुपर कंप्यूटर अकादमी, NVIDIA ग्राफिक्स त्वरक का उपयोग करते हुए उच्च प्रदर्शन क्लस्टर कम्प्यूटिंग
[[] IXBT.com: NVIDIA CUDA - GPU पर गैर-ग्राफिक कंप्यूटिंग
[९] cgm.computergraphics.ru: CUDA प्रौद्योगिकी का परिचय
[१०] THG.ru: एनवीडिया क्यूडा: एक ग्राफिक्स कार्ड या सीपीयू मृत्यु पर कंप्यूटिंग?
[११] step3d.narod.ru: CUDA की मूल बातें, CUDA प्रोग्रामिंग (भाग २)
[१२] CUDA कंपाइलर ड्राइवर (NVCC)
[१३] क्यूडा में इनलाइन पीटीएक्स असेंबली का उपयोग करना
[१४] पीटीएक्स: समानांतर धागा निष्पादन आईएसए संस्करण ३.०
[१५] सीयूडीए एपीआई संदर्भ मैनुअल ( पीडीएफ , एचटीएमएल ऑनलाइन )

CUDA: ब्लॉक सिंक्रोनाइज़ेशन