कैसे एक __global__ समारोह मान या सी/सी की तरह बाहर तोड़ सकते हैं ++

Question

करता है हाल ही में मैं CUDA पर नौकरियों की तुलना स्ट्रिंग कर रहा हूँ, और मैं आश्चर्य जब यह सटीक स्ट्रिंग पाता है कि कैसे एक __global__ समारोह एक मूल्य लौट सकते हैं कि मैं मैं देख रहा हूँ

मेरा मतलब है, मुझे __global__ फ़ंक्शन की आवश्यकता है जिसमें एक बड़ी बड़ी स्ट्रिंग-पूल के बीच एक निश्चित स्ट्रिंग को खोजने के लिए बड़ी मात्रा में धागे हैं, और मुझे आशा है कि एक बार सटीक स्ट्रिंग पकड़े जाने के बाद, __global__ फ़ंक्शन सभी को रोक सकता है धागे और मुख्य समारोह में वापस लौटें, और मुझे बताता है "उसने ऐसा किया"!

मैं CUDA सी उपयोग कर रहा हूँ मैं संभवतः इस लक्ष्य को हासिल कर सकते हैं कैसे?

Answer 1

सभी थ्रेड के निष्पादन को बाधित करने के लिए एक थ्रेड के लिए CUDA (या NVIDIA GPUs पर) में कोई रास्ता नहीं है। जैसे ही परिणाम मिलते हैं, आप कर्नेल के तुरंत बाहर नहीं निकल सकते हैं, आज भी यह संभव नहीं है।

लेकिन एक थ्रेड परिणाम मिलने के बाद आप जितनी जल्दी हो सके से बाहर निकल सकते हैं। यहां एक मॉडल है कि आप यह कैसे करेंगे।

__global___ void kernel(volatile bool *found, ...) 
{ 
    while (!(*found) && workLeftToDo()) { 

     bool iFoundIt = do_some_work(...); // see notes below 

     if (iFoundIt) *found = true; 
    } 
} 

इस पर कुछ नोट्स।

1. volatile के उपयोग पर ध्यान दें। यह महत्वपूर्ण है।
2. सुनिश्चित करें कि आप found — प्रारंभ करें जो कि कर्नेल लॉन्च करने से पहले डिवाइस सूचक — से false होना चाहिए!
3. धागे तुरन्त बाहर निकलने नहीं होगा जब एक और धागा अपडेट found। वे अगली बार जब वे लूप के शीर्ष पर वापस आ जाएंगे तो वे बाहर निकलेंगे।
4. आप do_some_work मामलों को कैसे कार्यान्वित करते हैं। यदि यह बहुत अधिक काम (या बहुत परिवर्तनीय) है, तो परिणाम मिलने के बाद बाहर निकलने में देरी लंबी (या परिवर्तनीय) होगी।यदि यह बहुत कम काम है, तो आपके थ्रेड अपने काम का अधिकतर समय उपयोगी काम करने के बजाय found की जांच करेंगे।
5. do_some_work कार्य आवंटित करने के लिए भी जिम्मेदार है (यानी कंप्यूटिंग/सूचकांक में वृद्धि), और आप यह कैसे करते हैं, यह समस्या विशिष्ट है।
6. यदि आपके द्वारा लॉन्च किए गए ब्लॉक की संख्या मौजूदा जीपीयू पर कर्नेल की अधिकतम अधिभोग से काफी बड़ी है, और थ्रेड ब्लॉक के पहले चल रहे "लहर" में कोई मिलान नहीं मिलता है, तो यह कर्नेल (और नीचे दिया गया एक) डेडलॉक कर सकते हैं। यदि पहली लहर में एक मैच मिलता है, तो बाद के ब्लॉक केवल found == true के बाद चलेंगे, जिसका अर्थ है कि वे लॉन्च करेंगे, फिर तुरंत बाहर निकलें। समाधान केवल एक ही ब्लॉक को लॉन्च करना है जैसे कि निवासी एक साथ हो सकता है (उर्फ "अधिकतम लॉन्च"), और तदनुसार अपने कार्य आवंटन को अपडेट करें।
7. यदि कार्यों की संख्या अपेक्षाकृत छोटी है, तो आप while को if के साथ प्रतिस्थापित कर सकते हैं और कार्यों की संख्या को कवर करने के लिए पर्याप्त थ्रेड चला सकते हैं। फिर डेडलॉक के लिए कोई मौका नहीं है (लेकिन पिछले बिंदु का पहला भाग लागू होता है)।
8. workLeftToDo() समस्या-विशिष्ट है, लेकिन ऐसा करने के लिए कोई काम नहीं होने पर यह झूठ वापस आ जाएगा, ताकि हम पर कोई मेल नहीं खा सकें कोई मिलान नहीं मिला है।

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

__global___ void kernel(volatile bool *found, ...) 
{ 
    volatile __shared__ bool someoneFoundIt; 

    // initialize shared status 
    if (threadIdx.x == 0) someoneFoundIt = *found; 
    __syncthreads(); 

    while(!someoneFoundIt && workLeftToDo()) { 

     bool iFoundIt = do_some_work(...); 

     // if I found it, tell everyone they can exit 
     if (iFoundIt) { someoneFoundIt = true; *found = true; } 

     // if someone in another block found it, tell 
     // everyone in my block they can exit 
     if (threadIdx.x == 0 && *found) someoneFoundIt = true; 

     __syncthreads(); 
    } 
} 

इस तरह, ब्लॉक चुनावों प्रति एक धागा वैश्विक चर, और केवल धागे कि एक मैच कभी यह करने के बारे में लगता है, इसलिए वैश्विक स्मृति यातायात कम से कम है।

इसके अलावा: __global__ फ़ंक्शंस शून्य हैं क्योंकि यह परिभाषित करना मुश्किल है कि धागे के 1000s से एकल CPU थ्रेड में मूल्य कैसे वापस लाया जाए। उपयोगकर्ता के लिए डिवाइस में रिटर्न सरणी या शून्य-प्रतिलिपि मेमोरी का विरोध करने के लिए यह छोटा है जो उसके उद्देश्य के अनुरूप है, लेकिन सामान्य तंत्र बनाना मुश्किल है।

अस्वीकरण: ब्राउज़र में लिखा कोड, अवांछित, असत्यापित।

Answer 2

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

आपके पास समस्या के प्रकार के साथ प्रत्येक ब्लॉक में 1 डी ग्रिड के साथ 1 डी ग्रिड के अलावा कुछ भी उपयोग करना आवश्यक नहीं है क्योंकि स्ट्रिंग पूल वास्तव में अन्य समस्याओं की तरह 2 डी में विभाजित नहीं होता है (उदाहरण के लिए मैट्रिक्स गुणा)

मैं स्ट्रिंग पूल में 100 स्ट्रिंग्स का एक साधारण उदाहरण कहूंगा और आप चाहते हैं कि वे सभी को अनुक्रमिक रूप से समानांतर फैशन में चेक किया जाए।

//main 
//Should cudamalloc and cudacopy to device up before this code 
dim3 dimGrid(10, 1); // 1D grid with 10 blocks 
dim3 dimBlocks(10, 1); //1D Blocks with 10 threads 
fun<<<dimGrid, dimBlocks>>>(, Height) 
//cudaMemCpy answerIdx back to integer on host 

//kernel (Not positive on these types as my CUDA is very rusty 
__global__ void fun(char *strings[], char *stringToMatch, int *answerIdx) 
{ 
    int idx = blockIdx.x * 10 + threadIdx.x; 

    //Obviously use whatever function you've been using for string comparison 
    //I'm just using == for example's sake 
    if(strings[idx] == stringToMatch) 
    { 
     *answerIdx = idx 
    } 
} 

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

Answer 3

आप रोमांच महसूस करते हैं, गिरी निष्पादन को रोकने के लिए एक वैकल्पिक दृष्टिकोण सिर्फ निष्पादित करने के लिए

// (write result to memory here) 
__threadfence(); 
asm("trap;"); 

अगर एक जवाब पाया जाता है होगा।

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

यदि आप एक सुरक्षित और सरल समाधान की तलाश में हैं, तो इसके बजाय मार्क हैरिस के सुझाव के साथ जाएं।