क्या एएमडी ओपनसीएल कर्नेल में कंपाइलर के साथ लूप को अनलॉक करने का कोई तरीका है?

Question

मैं एएमडी और एनवीडिया जीपीयू के लिए ओपनसीएल के बीच प्रदर्शन अंतर का आकलन करने की कोशिश कर रहा हूं। मेरे पास एक कर्नेल है जो मैट्रिक्स-वेक्टर गुणा करता है। मैं पलों पर दो अलग-अलग प्रणालियों पर कर्नेल चला रहा हूं, मेरे लैपटॉप में उबंटू 12.04 और क्यूडीए 4.0 (जिसमें ओपनसीएल लाइब्रेरी और हेडर शामिल हैं) के साथ एक एनवीडिया जीटी 525 मीटर है और दूसरा एक एएमडी राडेन एचडी 77070 के साथ एक डेस्कटॉप है जो उबंटू के साथ है 12.04 और नवीनतम उत्प्रेरक ड्राइवर।

कर्नेल में मेरे पास दो #pragma unroll कथन हैं जो Nvidia OpenCL कार्यान्वयन (~ 6x) के लिए एक बड़ी गति-अप उत्पन्न करते हैं। हालांकि एएमडी ओपनसीएल संस्करण किसी भी गति का उत्पादन नहीं करता है। एएमडी एपीपी कर्नेल विश्लेषक के साथ कर्नेल को देखते हुए त्रुटि देता है कि अनलॉक का उपयोग नहीं किया जाता है क्योंकि यात्रा गणना ज्ञात नहीं है। तो मेरा सवाल है, #pragma unroll एएमडी ओपनसीएल के साथ काम करता है या क्या कोई विकल्प है (शायद एक कंपाइलर ध्वज जिसे मैं अनजान हूं)। मैं नीचे

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n) 
{ 
    float sum = 0.0f; 
    __global float* A; 
    int i; 
    int j = 0; 
    int indx = get_global_id(0); 
    __local float xs[12000]; 
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) { 
     xs[i] = x[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    A = &a[indx]; 
#pragma unroll 256 
    for(i = 0; i < n; i++) { 
     sum += xs[i] * A[j]; 
     j += m; 
    } 
    y[indx] = sum; 
} 

यह वही गिरी दोनों कार्यान्वयन में सही परिणाम का उत्पादन लेकिन #pragma उतारना आदेशों एएमडी के लिए कुछ भी नहीं करते गिरी (उन्हें बाहर टिप्पणी द्वारा जाँच) शामिल किया है।

Answer 1

यह दस्तावेज नहीं है, लेकिन यह वास्तव में #pragma unroll के साथ काम करना चाहिए। क्या आप यह देखने के लिए कंपाइलर लॉग देख सकते हैं कि अनलोल लागू किया गया है या नहीं? मुझे यकीन नहीं है कि क्या कर्नेल विश्लेषक ओपनसीएल रनटाइम के समान कंपेलर का उपयोग करता है, तो आप जांचना चाहेंगे।

अन्यथा, यदि आप जानते हैं कि n 256 के हिस्सों में आता है, तो आप 256 तत्वों के ब्लॉक पर एक लूप और 256 के निश्चित आकार के साथ एक दूसरे को मैन्युअल रूप से अनलोल कर सकते हैं, जो अनलॉल करना आसान हो सकता है। यह निश्चित रूप से समस्या को हल करेगा कि यात्रा गणना स्थिर रूप से ज्ञात नहीं है।

हालांकि, ध्यान रखें कि एक लूप को अनलॉक करना आमतौर पर जीत का बड़ा नहीं होता है, क्योंकि आपके पास गणना करने के लिए कई रजिस्ट्रार नहीं हैं। लूप अनोलिंग से बढ़े हुए रजिस्टर दबाव से स्पिलिंग पंजीकृत हो सकती है, जो धीमी है। आपको यह जांचना चाहिए कि एएमडी कार्ड पर वास्तव में कर्नेल कितनी तेजी से है। एक नया एनवीआईडीआईए ओपनसीएल कंपाइलर अनोल प्रोगमा से भी लाभ नहीं उठा सकता है।