कैसे सीयूडीए निरंतर स्मृति आवंटन काम करता है?

Question

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

मान लें कि .cu फ़ाइल में कई कर्नेल हैं, प्रत्येक 64K निरंतर स्मृति से कम उपयोग करते हैं। लेकिन कुल निरंतर स्मृति उपयोग 64K से अधिक है। क्या क्रमशः इन कर्नल को कॉल करना संभव है? क्या होता है यदि उन्हें अलग-अलग धाराओं का उपयोग करके समवर्ती रूप से बुलाया जाता है?

क्या होता है यदि बड़ी सीयूडीए गतिशील लाइब्रेरी होती है जिसमें निरंतर मेमोरी की अलग-अलग मात्रा का उपयोग करके कई कर्नेल होते हैं?

क्या होता है यदि दो अनुप्रयोग होते हैं जिनमें प्रत्येक उपलब्ध निरंतर स्मृति के आधे से अधिक की आवश्यकता होती है? पहला एप्लिकेशन ठीक चलाता है, लेकिन दूसरा ऐप कब विफल होगा? ऐप शुरू होने पर, cudaMemcpyToSymbol() कॉल या कर्नेल निष्पादन पर?

Answer 1

Parallel Thread Execution ISA Version 3.1 सेक्शन 5.1.3 निरंतर बैंकों पर चर्चा करता है।

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

एक साधारण कार्यक्रम निरंतर स्मृति के उपयोग को वर्णन करने के लिए इस्तेमाल किया जा सकता है।

__constant__ int kd_p1; 
__constant__ short kd_p2; 
__constant__ char kd_p3; 
__constant__ double kd_p4; 

__constant__ float kd_floats[8]; 

__global__ void parameters(int p1, short p2, char p3, double p4, int* pp1, short* pp2, char* pp3,  double* pp4) 
{ 
    *pp1 = p1; 
    *pp2 = p2; 
    *pp3 = p3; 
    *pp4 = p4; 
    return; 
} 

__global__ void constants(int* pp1, short* pp2, char* pp3, double* pp4) 
{ 
    *pp1 = kd_p1; 
    *pp2 = kd_p2; 
    *pp3 = kd_p3; 
    *pp4 = kd_p4; 
    return; 
} 

compute_30, sm_30 के लिए इस संकलित करें और cuobjdump -sass <executable or obj> पर अमल करने के लिए एकत्रित न

Fatbin elf code: 
================ 
arch = sm_30 
code version = [1,6] 
producer = cuda 
host = windows 
compile_size = 32bit 
identifier = c:/dev/constant_banks/kernel.cu 

    code for sm_30 
      Function : _Z10parametersiscdPiPsPcPd 
    /*0008*/  /*0x10005de428004001*/  MOV R1, c [0x0] [0x44];  // stack pointer 
    /*0010*/  /*0x40001de428004005*/  MOV R0, c [0x0] [0x150];  // pp1 
    /*0018*/  /*0x50009de428004005*/  MOV R2, c [0x0] [0x154];  // pp2 
    /*0020*/  /*0x0001dde428004005*/  MOV R7, c [0x0] [0x140];  // p1 
    /*0028*/  /*0x13f0dc4614000005*/  LDC.U16 R3, c [0x0] [0x144]; // p2 
    /*0030*/  /*0x60011de428004005*/  MOV R4, c [0x0] [0x158];  // pp3 
    /*0038*/  /*0x70019de428004005*/  MOV R6, c [0x0] [0x15c];  // pp4 
    /*0048*/  /*0x20021de428004005*/  MOV R8, c [0x0] [0x148];  // p4 
    /*0050*/  /*0x30025de428004005*/  MOV R9, c [0x0] [0x14c];  // p4 
    /*0058*/  /*0x1bf15c0614000005*/  LDC.U8 R5, c [0x0] [0x146]; // p3 
    /*0060*/  /*0x0001dc8590000000*/  ST [R0], R7;     // *pp1 = p1 
    /*0068*/  /*0x0020dc4590000000*/  ST.U16 [R2], R3;    // *pp2 = p2 
    /*0070*/  /*0x00415c0590000000*/  ST.U8 [R4], R5;    // *pp3 = p3 
    /*0078*/  /*0x00621ca590000000*/  ST.64 [R6], R8;    // *pp4 = p4 
    /*0088*/  /*0x00001de780000000*/  EXIT; 
    /*0090*/  /*0xe0001de74003ffff*/  BRA 0x90; 
    /*0098*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a8*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b8*/  /*0x00001de440000000*/  NOP CC.T; 
      ........................................... 


      Function : _Z9constantsPiPsPcPd 
    /*0008*/  /*0x10005de428004001*/  MOV R1, c [0x0] [0x44];  // stack pointer 
    /*0010*/  /*0x00001de428004005*/  MOV R0, c [0x0] [0x140];  // p1 
    /*0018*/  /*0x10009de428004005*/  MOV R2, c [0x0] [0x144];  // p2 
    /*0020*/  /*0x0001dde428004c00*/  MOV R7, c [0x3] [0x0];  // kd_p1 
    /*0028*/  /*0x13f0dc4614000c00*/  LDC.U16 R3, c [0x3] [0x4]; // kd_p2 
    /*0030*/  /*0x20011de428004005*/  MOV R4, c [0x0] [0x148];  // p3 
    /*0038*/  /*0x30019de428004005*/  MOV R6, c [0x0] [0x14c];  // p4 
    /*0048*/  /*0x20021de428004c00*/  MOV R8, c [0x3] [0x8];  // kd_p4 
    /*0050*/  /*0x30025de428004c00*/  MOV R9, c [0x3] [0xc];  // kd_p4 
    /*0058*/  /*0x1bf15c0614000c00*/  LDC.U8 R5, c [0x3] [0x6];  // kd_p3 
    /*0060*/  /*0x0001dc8590000000*/  ST [R0], R7; 
    /*0068*/  /*0x0020dc4590000000*/  ST.U16 [R2], R3; 
    /*0070*/  /*0x00415c0590000000*/  ST.U8 [R4], R5; 
    /*0078*/  /*0x00621ca590000000*/  ST.64 [R6], R8; 
    /*0088*/  /*0x00001de780000000*/  EXIT; 
    /*0090*/  /*0xe0001de74003ffff*/  BRA 0x90; 
    /*0098*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a8*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b8*/  /*0x00001de440000000*/  NOP CC.T; 
      ..................................... 

मैं एस.ए.एस.एस. के अधिकार के लिए एनोटेट आप देखना चाहिए।

sm30 पर आप देख सकते हैं कि पैरामीटर ऑफसेट 0x140 से शुरू होने वाले निरंतर बैंक 0 में पास किए जाते हैं। आप cuobjdump --dump-elf <executable or obj> निष्पादित यदि आप अन्य रोचक लगातार जानकारी प्राप्त कर सकते

उपयोगकर्ता परिभाषित निरंतर चर निरंतर बैंक 3.

में परिभाषित कर रहे हैं।

32bit elf: abi=6, sm=30, flags = 0x1e011e 
Sections: 
Index Offset Size ES Align Type Flags Link  Info Name 
    1  34 142 0 1 STRTAB  0 0  0 .shstrtab 
    2 176 19b 0 1 STRTAB  0 0  0 .strtab 
    3 314  d0 10 4 SYMTAB  0 2  a .symtab 
    4 3e4  50 0 4 CUDA_INFO  0 3  b .nv.info._Z9constantsPiPsPcPd 
    5 434  30 0 4 CUDA_INFO  0 3  0 .nv.info 
    6 464  90 0 4 CUDA_INFO  0 3  a .nv.info._Z10parametersiscdPiPsPcPd 
    7 4f4 160 0 4 PROGBITS  2 0  a .nv.constant0._Z10parametersiscdPiPsPcPd 
    8 654 150 0 4 PROGBITS  2 0  b .nv.constant0._Z9constantsPiPsPcPd 
    9 7a8  30 0 8 PROGBITS  2 0  0 .nv.constant3 
    a 7d8  c0 0 4 PROGBITS  6 3 a00000b .text._Z10parametersiscdPiPsPcPd 
    b 898  c0 0 4 PROGBITS  6 3 a00000c .text._Z9constantsPiPsPcPd 

.section .strtab 

.section .shstrtab 

.section .symtab 
index  value  size  info other shndx name 
    0   0  0  0  0  0  (null) 
    1   0  0  3  0  a  .text._Z10parametersiscdPiPsPcPd 
    2   0  0  3  0  7  .nv.constant0._Z10parametersiscdPiPsPcPd 
    3   0  0  3  0  b  .text._Z9constantsPiPsPcPd 
    4   0  0  3  0  8  .nv.constant0._Z9constantsPiPsPcPd 
    5   0  0  3  0  9  .nv.constant3 
    6   0  4  1  0  9  kd_p1 
    7   4  2  1  0  9  kd_p2 
    8   6  1  1  0  9  kd_p3 
    9   8  8  1  0  9  kd_p4 
    10   16  32  1  0  9  kd_floats 
    11   0  192  12  10  a  _Z10parametersiscdPiPsPcPd 
    12   0  192  12  10  b  _Z9constantsPiPsPcPd 

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