ओपनसीएल एएमडी बनाम एनवीआईडीआईए प्रदर्शन

Question

मैंने एक साधारण कर्नेल लागू किया जो किसी प्रकार का संकल्प है। मैंने इसे एनवीआईडीआईए जीटी 240 पर मापा। ओडीसीएल पर लिखे जाने पर सीयूडीए पर लिखा गया और 100 एमएस में 70 एमएस लगा। ठीक है, मैंने सोचा, एनवीआईडीआईए कंपाइलर को सीयूडीए के लिए बेहतर अनुकूलित किया गया है (या मैं कुछ गलत कर रहा हूं)। मुझे इसे एएमडी जीपीयू पर चलाने की ज़रूरत है, इसलिए मैं एएमडी एपीपी एसडीके में माइग्रेट हुआ। बिल्कुल वही कर्नेल कोड।

मैंने दो परीक्षण किए और उनके परिणाम मेरे लिए निराश हो गए: एचडी 6670 पर 200 एमएस और एचडी 5850 पर 70 एमएस (जीटी 240 + सीयूडीए के लिए एक ही समय)। और मुझे इस तरह के अजीब व्यवहार के कारणों में बहुत दिलचस्पी है।

सभी परियोजनाओं को क्रमशः एनवीआईडीआईए और एएमडी की नमूना परियोजनाओं से सेटिंग्स का उपयोग करके वीएस -2010 पर बनाया गया था।

कृपया, मेरी पोस्ट को एनवीआईडीआईए विज्ञापन के रूप में नहीं मानें। मैं काफी समझता हूं कि एचडी 5850 जीटी 240 की तुलना में अधिक शक्तिशाली है। केवल एक चीज जिसे मैं जानना चाहता हूं वह है कि इस तरह का अंतर क्यों है और समस्या को कैसे ठीक किया जाए।

अद्यतन। नीचे कर्नेल कोड है जो बेस एक में समान रूप से आकार की टेम्पलेट छवियों को देखता है। आधार छवि के प्रत्येक पिक्सेल को टेम्पलेट्स में से एक की संभावित उत्पत्ति के रूप में माना जाता है और इसे एक अलग थ्रेड द्वारा संसाधित किया जाता है। कर्नेल बेस छवि और प्रत्येक टेम्पलेट के प्रत्येक पिक्सेल के आर, जी, बी मानों की तुलना करता है, और यदि कम से कम एक अंतर diff पैरामीटर से अधिक है, तो संबंधित पिक्सेल को बेजोड़ माना जाता है। यदि गैर-पिक्सेल पिक्सल की संख्या maxNonmatchQt से कम है तो संबंधित टेम्पलेट हिट हो गया है।

__constant int tOffset = 8196; // one template size in memory (in bytes) 
__kernel void matchImage6(__global unsigned char* image, // pointer to the base image 
      int imgWidth, // base image width 
      int imgHeight, // base image height 
      int imgPitch, // base image pitch (in bytes) 
      int imgBpp, // base image bytes (!) per pixel 
      __constant unsigned char* templates, // pointer to the array of templates 
      int tWidth, // templates width (the same for all) 
      int tHeight, // templates height (the same for all) 
      int tPitch, // templates pitch (in bytes, the same for all) 
      int tBpp, // templates bytes (!) per pixel (the same for all) 
      int diff, // max allowed difference of intensity 
      int maxNonmatchQt, // max number of nonmatched pixels 
      __global int* result, // results 
          ) { 
int x0 = (int)get_global_id(0); 
int y0 = (int)get_global_id(1); 
if(x0 + tWidth > imgWidth || y0 + tHeight > imgHeight) 
    return; 
int nonmatchQt[] = {0, 0, 0, 0, 0, 0}; 
for(int y = 0; y < tHeight; y++) { 
    int ind = y * tPitch; 
    int baseImgInd = (y0 + y) * imgPitch + x0 * imgBpp; 
    for(int x = 0; x < tWidth; x++) { 
     unsigned char c0 = image[baseImgInd]; 
     unsigned char c1 = image[baseImgInd + 1]; 
     unsigned char c2 = image[baseImgInd + 2]; 
     for(int i = 0; i < 6; i++) 
      if(abs(c0 - templates[i * tOffset + ind]) > diff || 
          abs(c1 - templates[i * tOffset + ind + 1]) > diff || 
          abs(c2 - templates[i * tOffset + ind + 2]) > diff) 
       nonmatchQt[i]++; 
     ind += tBpp; 
     baseImgInd += imgBpp; 
    } 
    if(nonmatchQt[0] > maxNonmatchQt && nonmatchQt[1] > maxNonmatchQt && nonmatchQt[2] > maxNonmatchQt && nonmatchQt[3] > maxNonmatchQt && nonmatchQt[4] > maxNonmatchQt && nonmatchQt[5] > maxNonmatchQt) 
     return; 
} 
for(int i = 0; i < 6; i++) 
    if(nonmatchQt[i] < maxNonmatchQt) { 
     unsigned int pos = atom_inc(&result[0]) * 3; 
     result[pos + 1] = i; 
     result[pos + 2] = x0; 
     result[pos + 3] = y0; 
    } 
} 

कर्नेल रन विन्यास: AMD के लिए और (32, 16) NVIDIA के लिए वैश्विक काम आकार = (1900, 1200) स्थानीय काम आकार = (8 32,)।

निष्पादन समय: HD 5850 - 69 एमएस, HD 6670 - 200 एमएस, GT 240 - 100 ms।

मेरे कोड के बारे में कोई भी टिप्पणी भी अत्यधिक सराहना की जाती है।

Answer 1

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

प्रत्येक हार्डवेयर अपना स्वयं का प्रदर्शन बढ़ावा देता है, उदाहरण के लिए एनवीआईडीआईए से देशी_। तो आपको उस हार्डवेयर के बारे में अधिक जानने की आवश्यकता है जिस पर आप काम कर रहे हैं, जो वास्तव में काम कर सकता है। लेकिन मैं व्यक्तिगत रूप से अनुशंसा करता हूं कि ऐसे हार्डवेयर विशिष्ट अनुकूलन का उपयोग न करें जो आपके कोड की लचीलापन को प्रभावित कर सकता है।

आप दिखाए गए कुछ कागजात भी देख सकते हैं, सीयूडीए प्रदर्शन उसी एनवीआईडीआईए हार्डवेयर पर ओपनसीएल प्रदर्शन से काफी बेहतर है।

तो यह कोड लिखने के लिए हमेशा बेहतर होता है जो डिवाइस विशिष्ट अनुकूलन के बजाए अच्छी लचीलापन प्रदान करता है।

Answer 2

निष्पादन समय में अंतर संकलक के कारण होता है। आपका कोड आसानी से वेक्टरकृत किया जा सकता है। वेक्टर प्रकार char4 के सरणी के रूप में छवि और टेम्पलेट्स पर विचार करें (प्रत्येक char4 वेक्टर के आगे समन्वय हमेशा 0 है)।इसके बजाय 3 स्मृति में लिखा है:

unsigned char c0 = image[baseImgInd]; 
unsigned char c1 = image[baseImgInd + 1]; 
unsigned char c2 = image[baseImgInd + 2]; 

उपयोग केवल एक:

unsigned char4 c = image[baseImgInd]; 

भारी अगर के बजाय

:

if(abs(c0 - templates[i * tOffset + ind]) > diff || 
       abs(c1 - templates[i * tOffset + ind + 1]) > diff || 
       abs(c2 - templates[i * tOffset + ind + 2]) > diff) 
     nonmatchQt[i]++; 

उपयोग तेजी से:

unsigned char4 t = templates[i * tOffset + ind]; 
    nonmatchQt[i] += any(abs_diff(c,t)>diff); 

इस प्रकार आप को बढ़ाने के प्रदर्शन आपके कोड का 3 गुना तक (यदि कंपि लेयर स्वयं कोड को सदिश नहीं करता है)। मुझे लगता है कि एएमडी ओपनसीएल कंपाइलर ऐसे वेक्टरेशन और अन्य अनुकूलन नहीं करता है। मेरे अनुभव से एनवीआईडीआईए जीपीयू पर ओपनसीएल आमतौर पर सीयूडीए से तेज किया जा सकता है, क्योंकि यह अधिक निम्न स्तर है।