सीयूडीए में किस प्रकार के चर रजिस्ट्रार का उपभोग करते हैं?

Question

__global__ void add(int *c, const int* a, const int* b) 
{ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    int offset = x + y * gridDim.x; 
    c[offset] = a[offset] + b[offset]; 
} 

उपरोक्त उदाहरण में, मुझे लगता है कि x, y, offset रजिस्टरों में जबकि

• NVCC -Xptxas -v देता 4 registers, 24+16 bytes smem

• प्रोफाइलर शो 4 रजिस्टरों

  सहेजे जाते हैं

• और सिर PTX फ़ाइल:

  .reg .u16 %rh<4>; 
  .reg .u32 %r<9>;  
  .reg .u64 %rd<10>; 
  .loc 15 21 0 
  
  $LDWbegin__Z3addPiPKiS1_: 
  .loc 15 26 0 
  

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

मुझे यकीन नहीं है कि मैं सही हूं या नहीं।

Answer 1

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

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

nvcc असेंबलर के रजिस्टर आवंटन व्यवहार को प्रभावित करने के कुछ तरीके प्रदान करता है। आपके पास __launch_bounds__ है जो कंपाइलर को हेरिस्टिक संकेत प्रदान करता है जो रजिस्टर आवंटन को प्रभावित कर सकता है, और कंपाइलर/असेंबलर -maxrregcount तर्क (स्थानीय मेमोरी में स्पिलिंग के संभावित खर्च पर, जो प्रदर्शन को कम कर सकता है) लेता है। अस्थिर कीवर्ड nvopen64 आधारित कंपाइलर के पुराने संस्करणों में अंतर लाने के लिए प्रयोग किया जाता है और स्थानीय मेमोरी स्पिल व्यवहार को प्रभावित कर सकता है। लेकिन आप मूल सी कोड या पीटीएक्स असेंबली भाषा कोड में मनमाने ढंग से नियंत्रण आवंटन को नियंत्रित या संचालित नहीं कर सकते हैं।