मल्टी-जीपीयू क्यूडा गणना

Question

मैं बहु-जीपीयू प्रोग्रामिंग में नौसिखिया हूं और मेरे पास बहु-जीपीयू कंप्यूटिंग के बारे में कुछ प्रश्न हैं। उदाहरण के लिए चलो डॉट-उत्पाद उदाहरण लें। मैं एक सीपीयू-थ्रेड चला रहा हूं जो 2 बड़े सरणी ए [एन] और बी [एन] बनाता है। इन सरणी के आकार के कारण मुझे अपने डॉट उत्पाद की गणना 2 जीपीयू में विभाजित करने की जरूरत है, दोनों टेस्ला एम 2050 (गणना क्षमता 2.0)। समस्या यह है कि मुझे अपने सीपीयू-थ्रेड द्वारा नियंत्रित डू-लूप के अंदर कई बार इन डॉट-उत्पादों की गणना करने की आवश्यकता है। प्रत्येक डॉट-उत्पाद को पिछले के परिणाम की आवश्यकता होती है। मैंने 2 अलग-अलग धागे बनाने के बारे में पढ़ा है जो अलग-अलग 2 जीपीयू को अलग से नियंत्रित करते हैं (जैसा कि उदाहरण के अनुसार क्यूडा पर वर्णित है) लेकिन मुझे उनके बीच डेटा सिंक्रनाइज़ करने और उनका आदान-प्रदान करने के बारे में कोई जानकारी नहीं मिली है। क्या कोई दूसरा विकल्प है? मैं वास्तव में किसी भी तरह की मदद/उदाहरण की सराहना करता हूं। अग्रिम धन्यवाद!

Answer 1

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

तो, CUDA 4.0 से शुरू होने पर, आप आसानी से एकल-थ्रेडेड होस्ट प्रोग्राम से एकाधिक GPUs प्रोग्राम कर सकते हैं। Here are some slides I presented last year about this।

प्रोग्रामिंग कई GPUs के इस रूप में सरल किया जा सकता है:

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    kernel<<<blocks, threads>>>(args); 
} 

डॉट उत्पादों के अपने विशिष्ट उदाहरण के लिए, यदि आप एक प्रारंभिक बिंदु के रूप thrust::inner_product इस्तेमाल कर सकते हैं। मैं प्रोटोटाइप के लिए ऐसा करूँगा। लेकिन बैंडविड्थ बाधाओं के अंत में मेरी टिप्पणियां देखें।

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

// assume the deviceIDs of the two 2050s are dev0 and dev1. 
// assume that the whole vector for the dot product is on the host in h_data 
// assume that n is the number of elements in h_vecA and h_vecB. 

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
float result = 0.f; 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1); 
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1); 
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f); 
} 

(मैं स्वीकार करते हैं कि अनुक्रमण ऊपर सही यदि n numDevs के यहां तक ​​कि कई नहीं है नहीं है, लेकिन मैं तय है कि यह सरल है पाठक के लिए एक व्यायाम। :)

के रूप में छोड़ देंगे , और एक महान शुरुआत है। इसे पहले काम करें, फिर अनुकूलित करें।

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

आप कुछ इस तरह कर सकता है:

int bytes = sizeof(float) * n; 
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable); 
// ... then fill your input arrays h_vecA and h_vecB 


for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    cudaEventCreate(event[d])); 
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0); 
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0); 
    cudaHostGetDevicePointer(&dresults[d], results, 0); 
} 

... 

for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    int first = d * (n/d); 
    int last = (d+1)*(n/d)-1; 
    my_inner_product<<<grid, block>>>(&dresults[d], 
             vecA+first, 
             vecA+last, 
             vecB+first, 0.f); 
    cudaEventRecord(event[d], 0); 
} 

// wait for all devices 
float total = 0.0f; 
for (int d = 0; d < devs; d++) { 
    cudaEventSynchronize(event[d]); 
    total += results[numDevs]; 
} 

Answer 2

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

आप एमपीआई या कांटा का भी उपयोग कर सकते हैं, इस मामले में संचार नेटवर्क प्रोग्रामिंग के समान तरीके से संचार किया जा सकता है ... पाइप/सॉकेट या संचार (सिंक्रनाइज़ेशन) (ब्लॉकिंग) भेजता है और प्राप्त करता है।