CUDA

2013-04-20 13 views
9

के साथ स्थिरांक का उपयोग करना CUDA में स्थिरांक का उपयोग करने का सबसे अच्छा तरीका कौन सा है?CUDA

एक तरह से, निरंतर स्मृति में स्थिरांक परिभाषित करने के लिए है की तरह:

// CUDA global constants 
__constant__ int M; 

int main(void) 
{ 
    ... 
    cudaMemcpyToSymbol("M", &M, sizeof(M)); 
    ... 
} 

एक वैकल्पिक रास्ता सी पूर्वप्रक्रमक उपयोग करने के लिए होगा:

#define M = ... 

मैं सी पूर्वप्रक्रमक साथ स्थिरांक को परिभाषित करने लगता होगा बहुत तेज है। फिर एक CUDA डिवाइस पर निरंतर स्मृति का उपयोग करने के लाभ क्या हैं?

+1

संकलन समय पर ज्ञात स्थिरांक प्रीप्रोसेसर मैक्रोज़ (यानी # परिभाषित') का उपयोग करके परिभाषित किया जाना चाहिए। अन्य मामलों में, '__constant__' [चर] [http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#constant) एक विकल्प हो सकता है जो CUDA प्रोग्रामर कोड को अनुकूलित करने के लिए उपयोग करता है जो एक्सेस करता है गणना किए गए चर जो बदलते नहीं हैं। ध्यान दें कि प्रतीक का संदर्भ देने के लिए 'एम "का उपयोग अब cuda 5 में मान्य नहीं है। –

+0

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

+0

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

उत्तर

9
  1. स्थिरांक कि संकलन समय पर जाना जाता है पूर्वप्रक्रमक मैक्रो (जैसे #define) या C/C++ const वैश्विक/फ़ाइल दायरे में चर के माध्यम से का उपयोग कर परिभाषित किया जाना चाहिए।
  2. __constant__ memory का उपयोग उन कार्यक्रमों के लिए फायदेमंद हो सकता है जो कुछ मानों का उपयोग करते हैं जो कर्नेल की अवधि के लिए नहीं बदलते हैं और जिसके लिए कुछ एक्सेस पैटर्न मौजूद होते हैं (उदाहरण के लिए सभी धागे एक ही समय में एक ही मूल्य तक पहुंचते हैं)। यह स्थिरांक से बेहतर या तेज़ नहीं है जो ऊपर आइटम 1 की आवश्यकताओं को पूरा करता है।
  3. विकल्पों की संख्या एक कार्यक्रम के द्वारा किए जाने के लिए तो संख्या में अपेक्षाकृत छोटे हैं, और इन विकल्पों गिरी निष्पादन प्रभावित करते हैं, अतिरिक्त संकलन समय अनुकूलन के लिए एक संभावित दृष्टिकोण का उपयोग किया जाएगा templated code/kernels
+2

सी/सी ++ स्टाइल स्थिरांक, प्रीप्रोसेसर मैक्रोज़, या सी ++ टेम्पलेट्स का उपयोग कई कारणों से __constant__ मेमोरी का उपयोग करने से तेज़ हो सकता है: 1. कंपाइलर अतिरिक्त अनुकूलन लागू कर सकता है, और 2. निरंतर निर्देश को तत्काल के रूप में एम्बेड किया जा सकता है। निरंतर कैश एक्सेस अतिरिक्त विलंब जोड़ने वाले निरंतर कैश को याद कर सकते हैं। –

+0

प्रत्येक थ्रेड में प्रत्येक थ्रेड में कन्स्ट्रक्टर को कॉल की वजह से गैर-तुच्छ प्रकारों के 'डीडी स्थिरांक को परिभाषित नहीं किया जाएगा? –

4

नियमित सी/सी ++ शैली स्थिरांक: CUDA सी (स्वयं सी 99 का एक संशोधन) स्थिरांक पूर्ण संकलन समय इकाइयां हैं। यह शायद ही आश्चर्यजनक है कि एनवीसीसी में होने वाले अनुकूलन की मात्रा जीपीयू प्रसंस्करण की प्रकृति को बहुत ही शामिल है।

#define: मैक्रोज़ हमेशा बहुत ही सुरुचिपूर्ण लेकिन चुटकी में उपयोगी होते हैं।

__constant__ परिवर्तनीय विनिर्देशक, हालांकि मेरी राय में एक पूरी तरह से नया जानवर और कुछ गलत नामक है। मैं क्या एनवीडिया नीचे अंतरिक्ष में here है नीचे रख देगा:

__constant__ क्वालीफायर, वैकल्पिक रूप से __device__ के साथ उपयोग किया, वाणी एक चर:

निरंतर स्मृति अंतरिक्ष में
  • रहता है,
  • एक आवेदन का जीवनकाल है,
  • ग्रिड के भीतर और मेजबान से रनटाइम लाइब्रेरी (cudaGetSymbolAddress()/के माध्यम से सभी धागे से सुलभ है 10 cudaGetSymbolSize()/cudaMemcpyToSymbol()/cudaMemcpyFromSymbol())।

एनवीडिया के दस्तावेज़ निर्दिष्ट करता है कि __constant__रजिस्टर स्तर गति (लगभग शून्य विलंबता) पर उपलब्ध है यह एक ही लगातार एक ताना के सभी सूत्र द्वारा पहुँचा जा रहा है प्रदान की है।

उन्हें CUDA कोड में वैश्विक दायरे में घोषित किया गया है। व्यक्तिगत (और वर्तमान में चल रहे) अनुभव के आधार पर अलग-अलग संकलन के समय आपको इस विनिर्देशक से सावधान रहना होगा, जैसे कि आपके CUDA कोड (.cu और।cuh फाइलें) सी-सी ++ कोड से सी-स्टाइल हेडर में रैपर फ़ंक्शन डालकर।

पारंपरिक "निरंतर" निर्दिष्ट चर के विपरीत हालांकि इन्हें रनवेटाइम पर होस्टटाइम कोड से प्रारंभ किया जाता है जो डिवाइस मेमोरी आवंटित करता है और आखिरकार कर्नेल लॉन्च करता है। मैं दोहराता हूं कि मैं वर्तमान में काम कर रहा कोड हूं जो प्रदर्शित करता है, इसे कर्नेल निष्पादन से पहले cudaMemcpyToSymbol() का उपयोग कर रनटाइम पर सेट किया जा सकता है।

वे कम से कम एल 1 कैश स्तर की गति के बारे में कहने के लिए काफी आसान हैं जो पहुंच के लिए गारंटीकृत है।