एक यूआरएल

Question

के लिए रजिस्टर का उपयोग करने के लिए सीयूडीए को मजबूर करना मेरे पास मेरे कर्नेल में कई अप्रयुक्त रजिस्टर्स हैं। मैं क्यूडीए को कुछ डेटा रखने के लिए कुछ रजिस्टरों का उपयोग करने के लिए बताना चाहता हूं, हर बार जब मुझे इसकी आवश्यकता होती है तो वैश्विक डेटा पढ़ने के बजाए। (मैं साझा मेम उपयोग करने में सक्षम नहीं हूँ।)

__global__ void simple(float *gData) { 
float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

संकलन w /: NVCC -arch sm_20 --ptxas-विकल्प = -v simple.cu, और मैं
0 बाइट्स फ्रेम ढेर मिलता है, 0 बाइट्स फैल स्टोर, 0 बाइट्स फैल भार
प्रयुक्त 2 रजिस्टर, 40 बाइट्स cmem [0]

__global__ void simple(float *gData) { 
register float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

रजिस्टर घोषणा कुछ नहीं करता है।
0 बाइट्स ढेर फ्रेम, 0 बाइट्स फैल स्टोर, 0 बाइट्स फैल भार
प्रयुक्त 2 रजिस्टर, 40 बाइट्स cmem [0]

__global__ void simple(float *gData) { 
volatile float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

अस्थिर घोषणा ढेर भंडारण बनाता है:
4096 बाइट्स फ्रेम ढेर, 0 बाइट्स फैल स्टोर, 0 बाइट्स फैल भार
प्रयुक्त 21 प्रतिरोधकारियों, 40 बाइट्स cmem [0]

1) एक आसान तरीका एक चर के लिए रजिस्टर अंतरिक्ष उपयोग करने के लिए संकलक बताने के लिए है?
2) 'स्टैक फ्रेम' कहां है: रजिस्टर, ग्लोबल मेम, स्थानीय मेम, ...? एक ढेर फ्रेम क्या है? (जब GPU एक ढेर एक आभासी ढेर है के बाद से?)
3) simple.ptx फ़ाइल मूल रूप से रिक्त है: (NVCC -arch sm_20 -ptx simple.cu)

.loc 2 14 2 
ret; 

किसी भी विचार है जहाँ मैं असली मशीन/संकलित कोड पा सकते हैं?

Answer 1

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

एसएम 2.0 जीपीयू (फर्मि) केवल थ्रेड के 63 रजिस्टरों तक का समर्थन करता है। यदि यह पार हो गया है, तो रजिस्टर मान को स्थानीय (ऑफ-चिप) मेमोरी से कैश पदानुक्रम द्वारा समर्थित/भर दिया जाएगा। एसएम 3.5 जीपीयू इसे प्रति थ्रेड 255 रजिस्ट्रार तक बढ़ाते हैं।

सामान्य रूप से, जैसा कि जेरेड उल्लेख करता है, प्रति थ्रेड के बहुत से रजिस्टरों का उपयोग करना वांछनीय नहीं है क्योंकि यह अधिभोग को कम करता है, और इसलिए कर्नेल में विलंबता छिपाने की क्षमता को कम करता है। जीपीयू समांतरता पर बढ़ते हैं और अन्य धागे से काम के साथ स्मृति विलंबता को कवर करके ऐसा करते हैं।

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

उदाहरण आप साझा स्मृति अगर के लिए एक मामला हो सकता है दे:

1. ब्लॉक में कई सूत्र एक ही डेटा का उपयोग करें, या
2. प्रति-धागा सरणी आकार काफी छोटा आवंटित करने के लिए है एकाधिक थ्रेड ब्लॉक में सभी धागे के लिए पर्याप्त स्थान (प्रति थ्रेड 1024 फ्लोट बहुत अधिक है)।

njuffa उल्लेख किया है, कारण अपने कर्नेल केवल 2 रजिस्टरों का उपयोग करता है क्योंकि आप कुछ भी कर्नेल में डेटा के साथ उपयोगी नहीं करते है, और मृत कोड सभी संकलक द्वारा हटा दिया गया।

Answer 2

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

इसके अलावा, सीयूडीए 4.0 के लिए कमांड लाइन स्विच -Xopencc=-O3 का उपयोग करने के लिए कुछ भी है ताकि सादा स्केलर्स (जैसे डाटा स्ट्रक्चर) रजिस्टरों में मैप किए गए हों (this post देखें)। CUDA> 4.0 के लिए आपको डीबग समर्थन अक्षम करना होगा (-G कमांड लाइन विकल्प - ऑप्टिमाइज़ेशन तब होता है जब डिबगिंग अक्षम होती है)।

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

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