ध्यान दें कि वास्तव में दो प्रकार के हैं। आपके पास CL_MEM_READ_ONLY, CL_MEM_WRITE_ONLY और CL_MEM_READ_WRITE हैं जब आपके बफर आवंटित करते हैं लेकिन फिर आपके पास कर्नेल कोड में अपने पॉइंटर्स को सजाने के लिए __read_only, __write_only और __read_write भी है।
इन दोनों का उपयोग अनुकूलन और त्रुटि जांच दोनों के लिए किया जा सकता है। चलिए पहले प्रदर्शन को देखते हैं। यदि केवल लिखने वाले बफर का सामना करना पड़ता है, तो लिखने के लिए लिखने के लिए अधिक कैश को सहेजने के लिए लिखने की आवश्यकता नहीं होती है (जैसे कैश के माध्यम से लिखना)। यह GPU हार्डवेयर पर बहुत निर्भर करता है और कम से कम NVIDIA हार्डवेयर में वास्तव में इसे लागू करने के लिए आवश्यक निर्देश हैं (.cs और .lu संशोधक)। आप उनके PTX ISA का उल्लेख कर सकते हैं।
st.global.u32 [%r10], %r11; // no cache operation specified
यह समझ में CUDA तो संकलक उन क्वालिफायर के लिए समकक्ष नहीं है के रूप में करता है:
__kernel void Memset4(__global __write_only unsigned int *p_dest,
const unsigned int n_dword_num)
{
unsigned int i = get_global_id(0);
if(i < n_dword_num)
p_dest[i] = 0; // this
}
के रूप में संकलित हो जाता है: मैं, जैसे वास्तव में इस अनुकूलन प्रदर्शन संकलक का कोई सबूत नहीं देखा है सबसे अधिक संभावना है कि चुपचाप उनको अनदेखा करता है। लेकिन उन्हें वहां रखने में कोई दिक्कत नहीं होती है, हम भविष्य में भाग्यशाली हो सकते हैं। सीयूडीए में, इस कार्यक्षमता में से कुछ को फ़ंक्शन का उपयोग करके और एल 1 (-Xptxas -dlcm=cg) में वैश्विक मेमोरी स्थानांतरण को कैशिंग करने में ऑप्ट इन/आउट करने के लिए कंपाइलर झंडे का उपयोग करके उजागर किया गया है। यदि आप पाते हैं कि कैश को छोड़कर आप एक बड़ा फायदा उठाते हैं तो आप हमेशा asm का उपयोग भी कर सकते हैं।
त्रुटि जांच के लिए, केवल पढ़ने के लिए बफर को लिखना कर्नेल घोषणा में const विनिर्देशक का उपयोग करके आसानी से बचा जाता है। शुद्ध "सी" में केवल लिखने वाले बफर से पढ़ने की अनुमति नहीं है।
मेमोरी होस्ट करने के लिए उन बफरों को मैप करते समय एक और संभावित अनुकूलन होता है। CL_MEM_READ_ONLY बफर मैपिंग करते समय, मैप किए गए क्षेत्र को अनियमित किया जा सकता है क्योंकि होस्ट केवल उस स्मृति को लिखता है, डिवाइस को केवल इसे पढ़ने के लिए। इसी तरह, जब CL_MEM_WRITE_ONLY बफर को अन-मैपिंग करते हैं, तो ड्राइवर को होस्ट मेमोरी से डिवाइस मेमोरी में (होस्ट द्वारा संभावित रूप से संशोधित) की प्रतिलिपि बनाने की आवश्यकता नहीं होती है। मैंने इसे माप नहीं लिया।
एक तरफ ध्यान दें के रूप में, मैं का उपयोग कर की कोशिश की है:
inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src)
{
#ifdef NVIDIA
unsigned int n_result;
asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src));
return n_result;
#else // NVIDIA
return *p_src; // generic
#endif // NVIDIA
}
inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value)
{
#ifdef NVIDIA
asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory");
#else // NVIDIA
*p_dest = n_value; // generic
#endif // NVIDIA
}
जो तुम भी sm_35 उपकरणों के साथ एक सरल memcpy कर्नेल पर लगभग 15 अतिरिक्त जीबी/सेकंड देता है (GTX 780 और K40 पर परीक्षण)। sm_30 पर ध्यान देने योग्य स्पीडअप नहीं देखा है (सुनिश्चित नहीं है कि इसका यहां भी समर्थन किया जाना चाहिए - हालांकि निर्देशों को पीटीएक्स से अलग नहीं किया जा रहा है)। ध्यान दें कि आपको NVIDIA स्वयं को परिभाषित करने की आवश्यकता है (या Detect OpenCL device vendor in kernel code देखें)।