GPU SSSP कार्यान्वयन

अमूर्त

इस लेख में मैं आपको एसएसएसपी एल्गोरिदम को कुशलता से समानांतर करने का तरीका बताना चाहता हूं - ग्राफिक्स त्वरक का उपयोग करके ग्राफ़ में सबसे छोटा रास्ता खोजना। ग्राफिक्स त्वरक के रूप में, केपलर वास्तुकला के GTX टाइटन कार्ड पर विचार किया जाएगा।

परिचय

हाल ही में, ग्राफिक्स त्वरक (जीपीयू) गैर-ग्राफिकल कंप्यूटिंग में तेजी से महत्वपूर्ण भूमिका निभा रहे हैं। उनके उपयोग की आवश्यकता उनकी अपेक्षाकृत उच्च उत्पादकता और कम लागत के कारण है। जैसा कि आप जानते हैं, GPU पर, संरचनात्मक ग्रिड पर समस्याओं को अच्छी तरह से हल किया जाता है, जहां समानता आसानी से एक या दूसरे तरीके से प्रतिष्ठित होती है। लेकिन ऐसे कार्य हैं जिन्हें बड़ी क्षमताओं की आवश्यकता होती है और गैर-संरचनात्मक ग्रिड का उपयोग करते हैं। इस तरह की समस्या का एक उदाहरण है सिंगल शॉर्टेस्ट सोर्स पाथ प्रॉब्लम (SSSP) - किसी दिए गए वर्टेक्स से सबसे छोटे रास्तों को सभी भारित ग्राफ में खोजने का कार्य। सीपीयू पर इस समस्या को हल करने के लिए, कम से कम दो प्रसिद्ध एल्गोरिदम हैं: डेस्ट्रा एल्गोरिदम और फोर्ड-बेलमैन एल्गोरिथम। GPU पर Deystra और Ford-Bellman एल्गोरिथ्म के समानांतर कार्यान्वयन भी हैं। यहाँ मुख्य लेख हैं जो इस समस्या के समाधान का वर्णन करते हैं:

1. CUDA, पवन हरीश और PJ नारायणन का उपयोग करके GPU पर बड़े ग्राफ़ एल्गोरिदम को त्वरित करना
2. सबसे कम पथ समस्या के लिए एक नया जीपीयू-आधारित दृष्टिकोण, हेक्टर ऑर्टेगा-अरेंज, यूरी टॉरेस, डिएगो आर। ललनोस और आर्टुरो गोंजालेज-एस्क्रिबानो

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

डेटा संरचना विवरण

हम एक अप्रत्यक्ष भारित ग्राफ की भंडारण संरचना पर संक्षेप में विचार करते हैं, क्योंकि भविष्य में इसका उल्लेख और रूपांतरण किया जाएगा। ग्राफ संकुचित सीएसआर प्रारूप में निर्दिष्ट है। इस प्रारूप का उपयोग व्यापक रूप से विरल मैट्रिस और रेखांकन के भंडारण के लिए किया जाता है। एन कोने और एम किनारों के साथ एक ग्राफ के लिए, तीन सरणियों की आवश्यकता होती है: xadj, adjncy, भार। Xadj सरणी का आकार N + 1 है, अन्य दो 2 * M हैं, क्योंकि किसी भी जोड़ी के लिए एक गैर-उन्मुख ग्राफ में सीधा और रिवर्स आर्क्स स्टोर करना आवश्यक है।

ग्राफ को संचय करने का सिद्धांत इस प्रकार है। इंडेक्स xadj [I] से xadj [I + 1] तक वर्जन I के पड़ोसियों की पूरी सूची इसमें नहीं है। इसी तरह के सूचकांक वर्टेक्स I. से प्रत्येक किनारे के वजन को स्टोर करते हैं। उदाहरण के लिए, बाईं ओर का आंकड़ा आसन्न मैट्रिक्स का उपयोग करके लिखे गए 5 कोने का एक ग्राफ दिखाता है, और सीएसआर प्रारूप में दाईं ओर।

एल्गोरिथ्म का GPU कार्यान्वयन

इनपुट तैयार करना

एक स्ट्रीमिंग मल्टीप्रोसेसर (SMX) पर कम्प्यूटेशनल लोड बढ़ाने के लिए, इनपुट डेटा को परिवर्तित करना आवश्यक है। सभी परिवर्तनों को दो चरणों में विभाजित किया जा सकता है:

• समन्वय प्रारूप (सीओओ) के लिए सीएसआर प्रारूप का विस्तार
• COO प्रारूप को क्रमबद्ध करें

पहले चरण में, इस प्रकार सीएसआर प्रारूप का विस्तार करना आवश्यक है: हम एक और शुरुआत सरणी पेश करते हैं जिसमें हम आर्क की शुरुआत को संग्रहीत करेंगे। फिर adjncy array में उनके सिरे स्टोर हो जाएंगे। इस प्रकार, पड़ोसियों को संग्रहीत करने के बजाय, हम आर्क्स स्टोर करेंगे। ऊपर वर्णित ग्राफ़ पर इस परिवर्तन का एक उदाहरण:



दूसरे चरण में, प्राप्त किनारों को क्रमबद्ध करना आवश्यक है ताकि प्रत्येक जोड़ी (यू, वी) ठीक एक बार हो। इस प्रकार, एक किनारे (यू, वी) को संसाधित करते समय, GPU के वैश्विक मेमोरी से इस बढ़त के बारे में डेटा को फिर से पढ़ने के लिए बिना एक किनारे (वी, यू) को संसाधित करना संभव है।

कोर कंप्यूटिंग कोर

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

int k = blockIdx.x * blockDim.x + threadIdx.x; if(k < maxV) { double w = weights[k]; unsigned en = endV[k]; unsigned st = startV[k]; if(dist[st] > dist[en] + w) // (*) { dist[st] = dist[en] + w; modif[0] = iter; } else if(dist[en] > dist[st] + w) // (**) { dist[en] = dist[st] + w; modif[0] = iter; } } 

इस कर्नेल में, प्रत्येक थ्रेड दो किनारों (आगे और पीछे) को संसाधित करता है, उनमें से एक के साथ दूरी को बेहतर बनाने की कोशिश करता है। यह स्पष्ट है कि यदि एक साथ ब्लॉक की दोनों शर्तों को पूरा नहीं किया जा सकता है। फोर्ड-बेलमैन एल्गोरिथ्म के विपरीत, जहां प्रत्येक किनारे को क्रमिक रूप से स्कैन किया जाता है, GPU-कार्यान्वित एल्गोरिथ्म में, प्रवाह की "दौड़" की स्थिति उत्पन्न हो सकती है - जब दो या अधिक प्रवाह समान दूर [I] सेल को अपडेट करने के लिए आते हैं। हम दिखाते हैं कि इस मामले में एल्गोरिथ्म सही है।

मान लीजिए कि दो थ्रेड K1 और K2 हैं जो सेल डिस्ट [I] को अपडेट करना चाहते हैं। इसका मतलब है कि शर्त (*) या (**) पूरी हो गई है। दो मामले संभव हैं। पहले - दो धागों में से एक ने सबसे छोटा मान दर्ज किया। फिर इन दो थ्रेड्स के लिए अगले पुनरावृत्ति पर स्थिति गलत होगी, और सेल डिस्ट [I] में मूल्य न्यूनतम होगा। दूसरा - दो थ्रेड्स में से एक न्यूनतम मूल्य दर्ज नहीं किया गया। फिर, अगले पुनरावृत्ति पर, स्थिति थ्रेड्स में से एक के लिए सही होगी, और दूसरे के लिए झूठी होगी। इस प्रकार, परिणाम दोनों मामलों में समान होगा, लेकिन एक अलग संख्या में पुनरावृत्तियों में प्राप्त किया गया।

Ford-Bellman कलन विधि के अनुकूलित संस्करण के अनुसार, यदि किसी भी पुनरावृत्ति में विकृत सरणी में कोई परिवर्तन नहीं हुए थे, तो आगे पुनरावृति का कोई मतलब नहीं है। इन उद्देश्यों के लिए जीपीयू पर मॉडिफ़ चर को पेश किया गया था, जिसमें थ्रेड्स ने वर्तमान पुनरावृत्ति की संख्या लिखी थी।

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

कार्यान्वित एल्गोरिथ्म का अनुकूलन

अगला, हम कुछ अनुकूलन पर विचार करते हैं जो एल्गोरिथ्म के प्रदर्शन में काफी सुधार कर सकते हैं। अंतिम वास्तुकला का ज्ञान अनुकूलन करने में सहायक हो सकता है।

आधुनिक सीपीयू में तीन-स्तरीय कैश होता है। पहले स्तर के कैश का आकार 64K है और सभी प्रोसेसर कोर पर निहित है। दूसरे स्तर के कैश का आकार 1 से 2 एमबी तक भिन्न होता है। तीसरे स्तर का कैश पूरे सीपीयू के लिए सामान्य है और इसका आकार 12-15 एमबी है।

आधुनिक जीपीयू में दो-स्तरीय कैश है। पहले स्तर के कैश का आकार 64KB है। साझा मेमोरी और रजिस्टरों को बाहर निकालने के लिए उपयोग किया जाता है। साझा मेमोरी के लिए 48KB से अधिक उपलब्ध नहीं है। यह प्रत्येक कंप्यूटिंग यूनिट में निहित है। दूसरे स्तर के कैश का अधिकतम आकार 1.5 एमबी है और पूरे जीपीयू के लिए सामान्य है। GPU की वैश्विक मेमोरी से डाउनलोड किए गए डेटा को कैश करने के लिए उपयोग किया जाता है। सबसे आधुनिक GPU GK110 चिप में 15 प्रसंस्करण इकाइयाँ हैं। यह पता चला है कि पहले स्तर के कैश का लगभग 48KB और दूसरे स्तर के कैश का 102KB एक ब्लॉक पर आते हैं। सीपीयू की तुलना में, यह बहुत छोटा है, इसलिए जीपीयू की वैश्विक मेमोरी से रीड ऑपरेशन केंद्रीय प्रोसेसर की मुख्य मेमोरी से अधिक महंगे हैं। इसके अलावा, केपलर आर्किटेक्चर में रीड-ओनली टेक्सचर कैश को सीधे एक्सेस करने की क्षमता है। ऐसा करने के लिए, संबंधित औपचारिक पैरामीटर से पहले कर्नेल में const __restrict जोड़ें।

बनावट कैश का उपयोग करना

इस कार्य में, आपको दूरस्थ दूरी को लगातार अपडेट और पढ़ने की आवश्यकता है। यह सरणी जीपीआर की वैश्विक मेमोरी में आर्क और उनके वजन के बारे में जानकारी की तुलना में काफी कम जगह लेती है। उदाहरण के लिए, 2 ²⁰ (लगभग 1 मिलियन) के वर्टिस की संख्या वाले ग्राफ के लिए, डिस्टैंट अरेंज 8 MB होगा। इसके बावजूद, इस सरणी तक पहुंच यादृच्छिक रूप से की जाती है, जो GPU के लिए बुरा है, क्योंकि वैश्विक मेमोरी से प्रत्येक कंप्यूटिंग ताना के अतिरिक्त डाउनलोड उत्पन्न होते हैं। प्रति ताना डाउनलोड की संख्या को कम करने के लिए, आपको L2 कैश में डेटा को सहेजना होगा, पढ़ना होगा, लेकिन अन्य ताना डेटा द्वारा उपयोग नहीं किया जाएगा।

चूंकि बनावट कैश केवल-पढ़ने के लिए है, इसलिए इसका उपयोग करने के लिए, मुझे समान दूरी के समान सरणी में दो अलग-अलग लिंक दर्ज करने होंगे। प्रासंगिक कोड:

 __global__ void relax_ker (const double * __restrict dist, double *dist1, … …) { int k = blockIdx.x * blockDim.x + threadIdx.x + minV; if(k < maxV) { double w = weights[k]; unsigned en = endV[k]; unsigned st = startV[k]; if(dist[st] > dist[en] + w) { dist1[st] = dist[en] + w; modif[0] = iter; } else if(dist[en] > dist[st] + w) { dist1[en] = dist[st] + w; modif[0] = iter; } } } 

नतीजतन, यह पता चला कि कर्नेल के अंदर सभी रीड ऑपरेशन एक सरणी के साथ किए जाते हैं, और सभी दूसरे के साथ ऑपरेशन लिखते हैं। लेकिन dist और dist1 लिंक दोनों एक ही GPU मेमोरी लोकेशन की ओर इशारा करते हैं।

बेहतर कैश उपयोग के लिए डेटा स्थानीयकरण

ऊपर वर्णित अनुकूलन के सर्वश्रेष्ठ प्रदर्शन के लिए, यह आवश्यक है कि डाउनलोड किया गया डेटा L2 कैश में यथासंभव लंबे समय तक हो। डिस्ट सरणी को एंडीवी और स्टार्ट वी सरणियों में संग्रहीत पूर्वनिर्धारित इंडेक्स का उपयोग करके एक्सेस किया जाता है। कॉल को स्थानीय बनाने के लिए, हम डिस्ट सरणी को एक निश्चित लंबाई के खंडों में विभाजित करते हैं, उदाहरण के लिए, पी तत्व। चूंकि ग्राफ में एन वर्टिकल हैं, इसलिए हमें अलग-अलग सेगमेंट (एन / पी + 1) मिलते हैं। अगला, हम इन खंडों में किनारों को निम्नानुसार सॉर्ट करेंगे। पहले समूह में, हम उन किनारों को असाइन करते हैं जिनके सिरे शून्य खंड में आते हैं, और शुरुआत - पहले शून्य में, फिर पहले में आदि। दूसरे समूह में, हम उन किनारों को असाइन करते हैं जिनके छोर पहले खंड में आते हैं, और शुरुआत - पहले शून्य में, फिर पहले में आदि।

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

एल्गोरिदम परीक्षा परिणाम

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

कोने की संख्या 2 ^ एन	संख्या की संख्या	चाप की संख्या	MB में विकृत सरणी का आकार	एमबी में पसलियों और वजन के सरणियों का आकार
14	16 384	524288	0,125	4
15	32768	1048576	0,250	8
16	65536	2 097 152	0,500	16
17	131 072	4 194 304	1	32
18	262 144	8 388 608	2	64
19	524288	16 777 216	4	128
20	1048576	33 554 432	8	256
21	2 097 152	६। १० 108६४	16	512
22	4 194 304	134 217 728	32	1024
23	8 388 608	268 435 456	64	2048
24	16 777 216	536 870 912	128	4096

यह तालिका से देखा जा सकता है कि 2 ¹⁸ या अधिक की संख्या के साथ ग्राफ के लिए दूरी दूरी सरणी पूरी तरह से GPU पर L2 कैश में फिट नहीं होती है। निविडिया जीटीएक्स टाइटन जीपीयू पर परीक्षण किया गया था, जिसमें 14 एसएमयू 192 क्यूडा कोर (कुल 2688) और 3.4 गीगाहर्ट्ज़ और 8 एमबी कैश की आवृत्ति के साथ तीसरी पीढ़ी का इंटेल कोर आई 7 प्रोसेसर है। सीपीयू पर प्रदर्शन की तुलना करने के लिए, अनुकूलित डिक्जस्ट्रा एल्गोरिदम का उपयोग किया गया था। सीपीयू पर काम करने से पहले डेटा क्रमांकन के रूप में कोई अनुकूलन नहीं किया गया था। समय के बजाय, प्रदर्शन सूचक समय की प्रति सेकंड संसाधित आर्क्स की संख्या है। इस मामले में, ग्राफ में आर्क्स की संख्या द्वारा प्राप्त समय को विभाजित करना आवश्यक है। अंतिम परिणाम के रूप में, औसतन 32 अंक लिए गए। अधिकतम और न्यूनतम मूल्यों की गणना भी की गई थी।

संकलन के लिए, 13 वें संस्करण के इंटेल कंपाइलर और NVCC CUDA 5.5 के साथ झंडे -O3 –arch = sm_35 का उपयोग किया गया था।

किए गए कार्य के परिणामस्वरूप, निम्नलिखित ग्राफ पर विचार करें:



यह ग्राफ़ निम्न एल्गोरिदम के लिए औसत प्रदर्शन घटता दिखाता है:

1. सभी अनुकूलन के साथ कैश और& पर प्रतिबंधित - GPU एल्गोरिथ्म
2. कैश ऑफ - बेहतर कैशिंग के लिए क्रमचय का अनुकूलन किए बिना GPU एल्गोरिथ्म
3. बनावट कैश ऑप्टिमाइज़ेशन के बिना - GPU एल्गोरिदम को प्रतिबंधित करें
4. कैश और& बंद करें - अनुकूलन के बिना मूल GPU एल्गोरिथ्म
5. सीपीयू - सीपीयू पर बुनियादी एल्गोरिथ्म

यह देखा जा सकता है कि दोनों अनुकूलन प्रदर्शन को बहुत प्रभावित करते हैं। यह ध्यान देने योग्य है कि const __restrict अनुकूलन का दुरुपयोग प्रदर्शन को कम कर सकता है। परिणामी त्वरण इस ग्राफ में देखा जा सकता है:



ग्राफ़ दिखाता है कि, सीपीयू के विपरीत, जीपीयू में औसत से विचलन की एक बड़ी श्रृंखला है। यह प्रवाह की एक "दौड़" के रूप में एल्गोरिथ्म कार्यान्वयन की ख़ासियत के कारण है, शुरू से शुरू करने के बाद से हम पुनरावृत्तियों की एक अलग संख्या प्राप्त कर सकते हैं।

निष्कर्ष

किए गए कार्य के परिणामस्वरूप, SSSP एल्गोरिथ्म को विकसित, कार्यान्वित और अनुकूलित किया गया - एक ग्राफ़ में सबसे छोटे रास्तों की खोज। यह एल्गोरिथम किसी दिए गए शीर्ष से अन्य सभी में एक ग्राफ में सबसे छोटी दूरी की खोज करता है। जीटीएक्स टाइटन मेमोरी में फिट होने वाले सभी ग्राफ़ों में, अधिकतम प्रदर्शन को रेखांकन द्वारा 2 ²¹ तक की संख्या के साथ दिखाया गया है - लगभग 1100 मिलियन प्रति सेकंड। अधिकतम औसत त्वरण जो 1 से 4 मिलियन तक वर्टिकल की संख्या के साथ ग्राफ़ पर लगभग 40 था।

साहित्य

1. CUDA
2. राजा
3. फोर्ड बेलमैन एल्गोरिथम
4. दिज्क्स्त्र का मूल एल्गोरिथ्म