CUDA: __syncthreads() अगर बयान

मेरे पास CUDA सिंक्रनाइज़ करने के बारे में कोई प्रश्न है। विशेष रूप से, मुझे बयानों में सिंक्रनाइज़ करने के बारे में कुछ स्पष्टीकरण की आवश्यकता है। मेरा मतलब है, अगर मैं ब्लॉक के अंदर धागे के एक अंश द्वारा हिट किए गए एक कथन के दायरे में __syncthreads() डालता हूं, तो क्या होता है? मैंने सोचा कि कुछ थ्रेड "हमेशा के लिए" अन्य थ्रेड के लिए इंतजार करेंगे जो सिंक्रनाइज़िंग पॉइंट को नहीं दबाएंगे। तो, मैं ने लिखा है और निरीक्षण करने के लिए कुछ नमूना कोड निष्पादित:CUDA: __syncthreads() अगर बयान

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 

     index += gridSize; 
    } 
}

हैरानी की बात है पर्याप्त, मैंने देखा कि उत्पादन एक बहुत "सामान्य" था (64 तत्वों, 32 blocksize):

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

तो मैं संशोधित थोड़ा निम्नलिखित तरीके से मेरी कोड:

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 
      vett[index] = 3; 
     __syncthreads(); 

     index += gridSize; 
    } 
}

और उत्पादन किया गया था:

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3

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

स्रोत

2012-09-20 biagiop1986

अपने प्रश्न का उत्तर देने वाले लोगों को चेकमार्क (स्वीकृत उत्तर दें) दें। – Yakk

आपको __syncthreads() का उपयोग नहीं करना चाहिए जब तक कि एक थ्रेड ब्लॉक के भीतर सभी धागे में कथन तक पहुंच न जाए। programming guide (B.6) से:

__syncthreads() सशर्त कोड में अनुमति दी है, लेकिन केवल हूबहू संपूर्ण थ्रेड ब्लॉक भर में सशर्त मूल्यांकन करता है, अगर नहीं तो कोड निष्पादन लटका या अनायास ही साइड इफेक्ट का उत्पादन होने की संभावना है।

असल में, आपका कोड एक अच्छी तरह से निर्मित CUDA प्रोग्राम नहीं है।

स्रोत

2012-09-20 20:07:50

बेशक यह नहीं है! लेकिन मैंने इसे केवल अपने व्यवहार का निरीक्षण करने के उद्देश्य से लिखा था। – biagiop1986

@ biagiop1986: अच्छा ...आपके पास लाइब्रेरी कोड और हार्डवेयर का एक टुकड़ा है जो एक दस्तावेज के साथ आता है जो कहता है, "आपको एक्स नहीं करना चाहिए"। अब आप * हमें *, जनता से पूछ रहे हैं, यदि आप एक्स करते हैं तो क्या होता है - हम कैसे जानते हैं? विक्रेता से पूछो! क्या यह जानना पर्याप्त नहीं है कि कार्यक्रम खराब हो जाएगा? –

यह निर्भर करता है ... यह कहने का अधिकार है कि मुझे अपने कार्यक्रमों में इस तरह के कोड से बचना चाहिए क्योंकि यह खराब है (और, मैं कसम खाता हूं, मैं करूँगा!), लेकिन मैं 'कैसे' के बारे में उत्सुक था। और, इसके अलावा, मुझे अक्सर विक्रेता-स्पष्टीकरण की तुलना में समस्याओं के बारे में स्पष्टीकरण मिलते हैं। तो, मैं भविष्य में आपके पास हर कोडिंग समस्या के लिए हर किसी के बजाय आपको यहां वापस आऊंगा। Stackoverflow सबसे अच्छा है! वैसे, सभी को धन्यवाद। – biagiop1986

सीयूडीए मॉडल एमआईएमडी है लेकिन मौजूदा एनवीआईडीआईए जीपीयू __syncthreads() को थ्रेड के बजाय वार्प ग्रैन्युलरिटी पर लागू करते हैं। इसका मतलब है, ये warps inside a thread-block हैं जो सिंक्रनाइज़ नहीं हैं threads inside a thread-block। __syncthreds() बाधा को मारने या कार्यक्रम से बाहर निकलने के लिए थ्रेड-ब्लॉक के सभी 'warps' की प्रतीक्षा करता है। अधिक जानकारी के लिए Henry Wong's Demistifying paper देखें।

स्रोत

2012-09-20 20:27:17 ahmad

वह पेपर वास्तव में एक अच्छा संदर्भ है। मैं भूल गया था कि यह सशर्त शाखाओं को भी शामिल करता है। – tera

धन्यवाद, महान संसाधन। – biagiop1986

संक्षेप में, व्यवहार अपरिभाषित है। तो कभी-कभी ऐसा हो सकता है जो आप चाहते हैं, या ऐसा नहीं हो सकता है, या (काफी संभावना है) बस आपके कर्नेल को लटकाएगा या क्रैश करेगा।

तुम सच में कैसे चीजें आंतरिक रूप से काम करते हैं उत्सुक हैं, तो आप को याद है कि धागे स्वतंत्र रूप से निष्पादित नहीं की जरूरत है, लेकिन एक समय में एक ताना (32 धागे के समूह)।

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

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

सिंक्रनाइज़ेशन से पहले एक और सशर्त शाखा का सामना करना पड़ता है तो स्थिति अधिक जटिल हो जाती है। इस समस्या को उन पथों के ढेर के साथ हल किया जाता है जिन्हें अभी भी निष्पादित करने की आवश्यकता है (सौभाग्य से ढेर की वृद्धि सीमित है क्योंकि हमारे पास एक वार के लिए अधिकतम 32 अलग-अलग कोड पथ हो सकते हैं)।

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

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

आप इस विवरण से अपना रास्ता काम कर सकते हैं कि कोड का आपका विशेष भाग कैसा व्यवहार करता है। लेकिन ध्यान रखें कि यह सब अपरिभाषित है, इसकी कोई गारंटी नहीं है, और किसी विशिष्ट व्यवहार पर निर्भर किसी भी समय आपके कोड को तोड़ सकता है।

आप कुछ और विवरणों के लिए PTX manual पर विशेष रूप से bar.sync निर्देश के लिए देख सकते हैं कि __syncthreads() संकलित करता है। हेनरी वोंग का "Demystifying GPU Microarchitecture through Microbenchmarking" paper, अहमद द्वारा नीचे संदर्भित, पढ़ने के लायक भी है। भले ही पुरानी वास्तुकला और सीयूडीए संस्करण के लिए, सशर्त शाखाओं और __syncthreads() के बारे में अनुभाग अभी भी वैध रूप से मान्य होते हैं।

स्रोत

2012-09-20 20:28:03 tera

धन्यवाद, बहुत स्पष्ट व्याख्या। – biagiop1986

__syncthreads() का उपयोग ब्लॉक के भीतर धागे को सिंक्रनाइज़ करने के लिए किया जाता है। इसका मतलब है कि ब्लॉक में सभी धागे आगे बढ़ने से पहले पूरा होने की प्रतीक्षा करेंगे।

उस मामले पर विचार करें जहां ब्लॉक में कुछ धागे हैं, जो if-statement में प्रवेश करते हैं और कुछ नहीं करते हैं। इंतजार कर रहे धागे, अवरुद्ध किया जाएगा; हमेशा इंतज़ार कर रहा है।

आम तौर पर, यदि एक सशर्त कथन में सिंक्रनाइज़ करने के लिए यह एक अच्छी शैली नहीं है। इससे बचने के लिए सबसे अच्छा, और यदि आपके पास है तो अपने कोड को फिर से डिजाइन करें। सिंक्रनाइज़ करने का उद्देश्य यह सुनिश्चित करना है कि सभी धागे एक साथ आगे बढ़ें, आप पहली जगह में if-statement का उपयोग करके उन्हें क्यों फ़िल्टर करते हैं?

जोड़ने के लिए, यदि ब्लॉक में सिंक्रनाइज़ेशन की आवश्यकता है। कर्नेल का पुन: लॉन्च आवश्यक है।

स्रोत

2012-09-21 07:28:33

उत्तर

संबंधित मुद्दे