2012-08-03 15 views
5

मैं ओपनएसीसी (पीजीआई के कंपाइलर के साथ) सीख रहा हूं और मैट्रिक्स गुणा उदाहरण को अनुकूलित करने की कोशिश कर रहा हूं।ओपनएसीसी का उपयोग कर मैट्रिक्स गुणा को अनुकूलित करने के लिए कैसे?

void matrix_mul(float *restrict r, float *a, float *b, int N, int accelerate){ 

#pragma acc data copyin (a[0: N * N ], b[0: N * N]) copyout (r [0: N * N ]) if(accelerate) 
{ 
# pragma acc region if(accelerate) 
{ 
# pragma acc loop independent vector(32) 
for (int j = 0; j < N; j ++) 
{  
    # pragma acc loop independent vector(32) 
    for (int i = 0; i < N ; i ++) 
    { 
     float sum = 0; 
     for (int k = 0; k < N ; k ++) { 
     sum += a [ i + k*N ] * b [ k + j * N ]; 
     } 
     r[i + j * N ] = sum ; 
    } 
} 
} 
} 

यह आकार 32x32 धागे की धागा ब्लॉक में परिणाम है और मुझे सबसे अच्छा प्रदर्शन अब तक देता है: सबसे तेजी से कार्यान्वयन मैं अब तक आया निम्नलिखित है। यहाँ मानक हैं:

Matrix multiplication (1500x1500): 
GPU: Geforce GT650 M, 64-bit Linux 

Data sz    : 1500  
Unaccelerated: 
    matrix_mul() time : 5873.255333 msec 
Accelerated: 
    matrix_mul() time : 420.414700 msec 

Data size    : 1750 x 1750  
    matrix_mul() time : 876.271200 msec 
Data size    : 2000 x 2000  
    matrix_mul() time : 1147.783400 msec 
Data size    : 2250 x 2250  
    matrix_mul() time : 1863.458100 msec 
Data size    : 2500 x 2500  
    matrix_mul() time : 2516.493200 msec 

दुर्भाग्य से मुझे एहसास हुआ कि उत्पन्न CUDA कोड काफी आदिम है (उदाहरण के लिए तो यह और भी साझा स्मृति का उपयोग नहीं करता) और इसलिए हाथ से अनुकूलित CUDA कार्यक्रम के साथ प्रतिस्पर्धा नहीं कर सकते।

Arrayfire 1500 x 1500 matrix mul 
CUDA toolkit 4.2, driver 295.59 
GPU0 GeForce GT 650M, 2048 MB, Compute 3.0 (single,double) 
Memory Usage: 1932 MB free (2048 MB total) 
af: 0.03166 seconds 

Arrayfire 1750 x 1750 matrix mul 
af: 0.05042 seconds 
Arrayfire 2000 x 2000 matrix mul 
af: 0.07493 seconds 
Arrayfire 2250 x 2250 matrix mul 
af: 0.10786 seconds 
Arrayfire 2500 x 2500 matrix mul 
af: 0.14795 seconds 

मुझे आश्चर्य है कि अगर वहाँ किसी भी सुझाव कैसे OpenACC से बेहतर प्रदर्शन प्राप्त करने के लिए: एक संदर्भ कार्यान्वयन के रूप में मैं निम्नलिखित परिणामों के साथ Arrayfire lib ले गया? शायद निर्देशों की मेरी पसंद सही नहीं है?

+1

यह समस्या कंपाइलर निर्देश बनाम CUDA/OpenCL के विभिन्न दृष्टिकोण को दर्शाती है। कूडा/ओपनसीएल एच/डब्ल्यू के बहुत करीब है; जहां आप एक एच/डब्ल्यू मंच के लिए अनुकूलित और tweak कर सकते हैं। आप आंतरिक लूप कंप्यूटिंग 2,4, या 8 को अनलोल कर सकते हैं ... इस प्रकार आंतरिक लूपों की संख्या को कम कर देता है –

+1

हू अच्छा विचार, धन्यवाद .. हाँ मुझे पता है, CUDA/OpenCL को "निम्न-स्तरीय" एपीआई माना जा सकता है, मैं खुद पुराने कुडा स्कूल से हूं। दूसरी तरफ, ओपनएसीसी के पास भविष्य में अधिक क्षमता है क्योंकि यह केवल जीपीयू और निश्चित रूप से विकास लागत तक ही सीमित नहीं है। फिर भी, यह अच्छा होगा अगर ओपनएसीसी कंपाइलर्स कम्प्यूटेशंस के लिए जीपीयू की साझा मेमोरी का फायदा उठा सकें: मुझे पता है कि ओपनएसीसी 'कैश' निर्देश है लेकिन मैं इसे –

उत्तर

4

आप 14x स्पीडअप पर सही हो रहे हैं, जो मेरे अनुभव में पीजीआई के कंपाइलर के लिए बहुत अच्छा है।

सबसे पहले, क्या आप -Minfo के साथ संकलित कर रहे हैं? यह आपको अनुकूलन विकल्पों के संबंध में संकलक से बहुत सी प्रतिक्रिया देगा।

आप 32x32 थ्रेड ब्लॉक का उपयोग कर रहे हैं, लेकिन मेरे अनुभव में 16x16 थ्रेड ब्लॉक बेहतर प्रदर्शन प्राप्त करते हैं। यदि आप वेक्टर (32) खंडों को छोड़ देते हैं, तो कंपाइलर किस शेड्यूलिंग का चयन करता है?

प्रतिबंध के साथ ए और बी घोषित करने से संकलक बेहतर कोड उत्पन्न कर सकता है।

बस अपना कोड देखकर, मुझे यकीन नहीं है कि साझा स्मृति प्रदर्शन में मदद करेगी। साझा स्मृति केवल प्रदर्शन को बेहतर बनाने में मदद करती है यदि आपका कोड वैश्विक स्मृति पर जाने के बजाय वहां मूल्यों को संग्रहीत और पुन: उपयोग कर सकता है। इस मामले में आप इसे पढ़ने के बाद ए या बी के किसी भी भाग का पुन: उपयोग नहीं कर रहे हैं।

यह भी ध्यान देने योग्य है कि साझा स्मृति उपयोग की बात आने पर मुझे पीजीआई के कंपाइलर के साथ बुरे अनुभव हुए हैं। यह कभी-कभी मजाकिया सामान करता है और गलत मानों को कैश करता है (ऐसा लगता है कि यदि आप लूप को पिछड़ा करते हैं तो अधिकतर होता है), गलत परिणाम उत्पन्न करते हैं। मुझे वास्तव में साझा किए गए मेमोरी उपयोग को बाईपास करके, अनियंत्रित -ta = nvidia, nocache विकल्प को सही तरीके से काम करने के लिए अपने वर्तमान एप्लिकेशन को संकलित करना होगा।

+0

पर काम करने के लिए प्रबंधित नहीं कर सका, हाँ मैंने 16x16 मामले की कोशिश की लेकिन यह वास्तव में धीमी गति से चलता है। मुझे लगता है कि यह ठीक है क्योंकि कोई साझा स्मृति का उपयोग नहीं किया जाता है। इसलिए, प्रति ब्लॉक जितना अधिक धागे उतना अधिक होता है जितना अधिक "कैशिंग" इंटरमीडिएट परिणामों का परिणाम रजिस्टरों में होता है। वास्तव में एक तरीका है कि साझा स्मृति कैसे प्रदर्शन में मदद कर सकती है यदि आप CUDA SDK में मैट्रिक्स गुणा उदाहरण देखें। यदि मैं वेक्टर (32) क्लॉज को हटा देता हूं, तो कंपाइलर मैट्रिक्स की पंक्तियों (2 डी टाइल्स द्वारा नहीं) और प्रदर्शन बूंदों द्वारा बस vectorizes। वैसे भी एक अच्छी सलाह के लिए धन्यवाद –

संबंधित मुद्दे