CUDA: विशिष्ट कार्यों के प्रदर्शन पहलुओं

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

मान लें कि आपने लगभग यह पता लगा लिया है कि CUDA कैसे काम करता है और CUDA टूलकिट के स्थिर संस्करण को पहले ही डाउनलोड कर चुका है।

मैं अब कोर डुओ E8400 पर मिडिल-एंड GTX460 ग्राफिक्स कार्ड को पीड़ा देगा।

फंक्शन कॉल

हां, यदि हम कुछ गणना करना चाहते हैं, तो हम कार्ड पर किए गए फ़ंक्शन को कॉल किए बिना नहीं कर सकते। ऐसा करने के लिए, हम सबसे सरल परीक्षण फ़ंक्शन लिखते हैं:

__global__ void stubCUDA( unsigned short * output)
{
// the most valid function: yep, does nothing.
}

आपको याद दिला दूं कि __global__ Specier आपको सीपीयू से कॉल करके GPU पर एक फ़ंक्शन निष्पादित करने की अनुमति देता है:

cudaThreadSynchronize();
stubCUDA<<<GRID, THREADS>>>(0);
cudaThreadSynchronize();

सभी फ़ंक्शन कॉल डिफ़ॉल्ट रूप से अतुल्यकालिक होते हैं, इसलिए cudaThreadSynchronize () कॉल को कॉल करने के लिए पूर्ण होने की प्रतीक्षा करने की आवश्यकता होती है।

आइए इस तरह के ब्लॉक को लूप में चलाने का प्रयास करें: हमें GRID = 160, THREADS = 96 के लिए प्रति सेकंड लगभग 15,000 कॉल मिलते हैं।

मान लें कि यह बिलकुल भी मोटा नहीं है। यहां तक ​​कि सबसे सरल कार्य, जो कुछ भी नहीं करता है, 0.7 एमएस से अधिक तेजी से निष्पादित नहीं कर सकता है।

पहली धारणा यह है कि ज्यादातर समय थ्रेड सिंक्रोनाइज़ेशन पर खर्च किया जाता है और एसिंक्रोनस कॉल बहुत तेजी से काम करेंगे (हालांकि विशिष्ट कार्यों में उनका उपयोग करना अधिक विशिष्ट है)।

इसकी जाँच करें। सिंक्रनाइज़ेशन के बिना, फ़ंक्शन को प्रति सेकंड 73100 बार चलाना संभव था। परिणाम, यह ध्यान दिया जाना चाहिए, बिल्कुल प्रभावशाली नहीं है।

और अंतिम परीक्षण, चलो फ़ंक्शन को GRID = THREADS = 1 के साथ चलाते हैं, ऐसा लगता है कि यह कार्ड के अंदर थ्रेड का एक गुच्छा बनाने के ओवरहेड को समाप्त करना चाहिए। लेकिन ऐसा नहीं है, हमें प्रति सेकंड समान 73000-73500 कॉल मिलते हैं।

तो नैतिक:

• यह कार्ड पर चलने के लिए बिल्कुल व्यर्थ है जो सीपीयू पर मिलीसेकंड में गिने जाते हैं।
• एक कॉल के बाद थ्रेड सिंक्रोनाइज़ेशन, मध्यम-आकार के कार्यों पर प्रदर्शन को थोड़ा कम करता है।
• थ्रेड्स की संख्या और ग्रिड का आकार प्रति सेकंड कॉल की कुल संख्या को प्रभावित नहीं करता है (निश्चित रूप से, यह "उपयोगी" कार्यों के लिए ऐसा नहीं है जो कुछ करते हैं)।

बाहरी मेमोरी एक्सेस

कुछ उपयोगी खोजने के लिए, हमें इनपुट और आउटपुट डेटा की आवश्यकता होती है। ऐसा करने के लिए, आपको यह समझने की आवश्यकता है कि कितनी तेजी से डेटा को वीडियो कार्ड से / से स्थानांतरित किया जा रहा है। हम निम्नलिखित फ़ंक्शन का उपयोग करते हैं:

cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice);

हां, CUDA हमें अतुल्यकालिक डेटा ट्रांसमिशन टूल भी प्रदान करता है, लेकिन उनका प्रदर्शन, आगे देखते हुए, सिंक्रोनस फ़ंक्शन से भिन्न नहीं होता है।

हम बड़े ब्लॉक्स की नकल करते हैं: जैसे cudaMemcpyHostToDevice, और cudaMemcpyDeviceToHost की दिशा में हमें बड़े ब्लॉक (100 मेगाटन से अधिक) पर लगभग 2 GB / s का प्रदर्शन मिलता है। सामान्य तौर पर, यह बहुत अच्छा है।

बहुत छोटी संरचनाओं के साथ हालात बहुत खराब हैं। 4 बाइट्स ट्रांसमिट करके , हमें प्रति सेकंड 22,000 से अधिक कॉल नहीं मिलते हैं, अर्थात्। 88 केबी / एस ।

नैतिक:

• यह सलाह दी जाती है कि डेटा को बड़े ब्लॉकों में समूहित करें और इसे एक कॉल में cudaMemcpy फ़ंक्शन में पास करें।

आंतरिक मेमोरी एक्सेस

हमने डेटा को कार्ड में स्थानांतरित करने के बाद, आप उनके साथ काम करना शुरू कर सकते हैं। मैं वीडियो मेमोरी तक पहुंच की अनुमानित गति का मूल्यांकन करना चाहूंगा। ऐसा करने के लिए, हम निम्नलिखित फ़ंक्शन लिखते हैं:

__global__ void accessTestCUDA(unsigned short * output, unsigned short * data, int blockcount, int blocksize)
{
// just for test of max access speed: does nothing useful
unsigned short temp;
for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
temp = data[j];
}
}
output[0] = temp;
}


GRID और THREADS पैरामीटर पहले से ही यहां उपयोग किए जाते हैं, जब तक कि मैं क्यों समझाता हूं, लेकिन मेरा विश्वास करो - सब कुछ वैसा ही है जैसा कि यह होना चाहिए। पिकी कहेगा कि सिंक्रनाइज़ेशन की कमी के कारण परिणाम गलत तरीके से लिखा गया है, लेकिन हमें इसकी आवश्यकता नहीं है।

इसलिए, हमें मनमानी पढ़ने के लिए लगभग 42 जीबी / एस मिलता है। यह बिल्कुल भी बुरा नहीं है।

अब हम फ़ंक्शन को संशोधित करते हैं ताकि यह इनपुट डेटा को आउटपुट में कॉपी कर दे। इसका कोई मतलब नहीं है, लेकिन आपको वीडियो मेमोरी में रिकॉर्डिंग की गति का मूल्यांकन करने की अनुमति देता है (चूंकि परिवर्तन पूरी तरह से सरल है, मैं कोड की नकल नहीं करूंगा)।

I / O के लिए हमें लगभग 30 GB / s मिलते हैं। बुरा भी नहीं है।

यह सही होना चाहिए कि वास्तव में हमने अनुक्रमिक (कुछ विचलन के साथ) मेमोरी एक्सेस का उपयोग किया था। एक मनमाना आंकड़ा दो गुना तक बिगड़ सकता है - लेकिन क्या यह कोई समस्या नहीं है?

नैतिक:

• कार्ड पर मेमोरी तक पहुंच की बहुत उच्च गति के कारण, यह एल्गोरिदम को कुशलतापूर्वक लागू करना संभव है जो इसका गहन रूप से उपयोग करते हैं।

अंकगणित संचालन

हम बहुत ही सरल उदाहरणों को छोड़ते हैं और कुछ उपयोगी करते हैं। अर्थात्, छवि सामान्यीकरण (पिक्सेल [t]: = (पिक्सेल [t] -sub) * कारक)। वास्तव में कोड:

__global__ void normalizeCUDA(unsigned short * data, int blockcount, int blocksize, float sub, float factor)
{
for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
register float d = ( float )data[j];
d = (d - sub) * factor;
data[j] = (unsigned short )d;
}
}
}


यहां, तीन के रूप में कई उचित रूप से कम्प्यूटेशनल प्रक्रियाओं का उपयोग किया जाता है: वास्तविक संख्या, ADDMUL और पूर्णांक के लिए कास्टिंग। फ़ोरम डराता है कि पूरी सामग्री को कास्टिंग बहुत बुरी तरह से काम करता है। शायद यह पुरानी पीढ़ी के कार्ड के लिए सच था, लेकिन अब ऐसा नहीं है।

कुल प्रसंस्करण गति: 26 जीबी / एस । केवल 13% द्वारा I / O को निर्देशित करने के लिए तीन ऑपरेशनों ने प्रदर्शन को खराब कर दिया।

यदि आप ध्यान से कोड को देखते हैं, तो यह सामान्य रूप से सही नहीं है। पूर्णांक लिखने से पहले, वास्तविक को गोल होना चाहिए, उदाहरण के लिए गोल () फ़ंक्शन के साथ। लेकिन ऐसा न करें, और कभी भी इसका उपयोग करने की कोशिश न करें!

दौर (डी): 20 जीबी / एस , एक और शून्य से 23%।
(अहस्ताक्षरित लघु) (डी + 0.5): 26 जीबी / एस , वास्तव में माप त्रुटि के भीतर का समय भी नहीं बदला है।

नैतिक:

• अंकगणित संचालन वास्तव में तेजी से काम करते हैं!
• सरलतम इमेज प्रोसेसिंग एल्गोरिदम के लिए, आप 10-20 जीबी / एस की गति पर भरोसा कर सकते हैं।
• राउंड () फ़ंक्शन का उपयोग करने से बचना बेहतर है।

तार्किक संचालन

आइए तार्किक संचालन की गति का मूल्यांकन करने का प्रयास करें, और उसी समय हम एक और अच्छा काम करेंगे: हम सरणी में न्यूनतम और अधिकतम मान पाएंगे। यह चरण आम तौर पर सामान्यीकरण से पहले होता है (और यह सिर्फ इसके लिए लिखा गया था), लेकिन हमारे पास इसके विपरीत होगा - क्योंकि वह कठिन है। यहाँ काम कर कोड है:

__global__ void getMinMaxCUDA(unsigned short * output, unsigned short * data, int blockcount, int blocksize)
{
__shared__ unsigned short sMins[MAX_THREADS];
__shared__ unsigned short sMaxs[MAX_THREADS];

sMins[threadIdx.x] = data[0];
sMaxs[threadIdx.x] = data[0];

for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
register unsigned short d = data[j];
if (d < sMins[threadIdx.x])
sMins[threadIdx.x] = d;
if (d > sMaxs[threadIdx.x])
sMaxs[threadIdx.x] = d;
}
}

__syncthreads();

if (threadIdx.x == 0)
{
register unsigned short min = sMins[0];
for ( int j = 1; j < blockDim.x; j++)
if (sMins[j] < min)
min = sMins[j];
if (min < output[0])
output[0] = min;
}

if (threadIdx.x == 1)
{
register unsigned short max = sMaxs[0];
for ( int j = 1; j < blockDim.x; j++)
if (sMaxs[j] > max)
max = sMaxs[j];
if (max > output[1])
output[1] = max;
}

__syncthreads();
}



यहां आप थ्रेड्स और साझा की गई मेमोरी के सिंक्रनाइज़ेशन के बिना नहीं कर सकते।

कुल गति: 29 जीबी / एस , सामान्य से भी तेज।

मैंने न्यूनतम और अधिकतम कोड क्यों जोड़ा - दोनों की आमतौर पर आवश्यकता होती है, और कॉल अलग से समय खो देते हैं (पहले पैराग्राफ देखें)।

सामान्य तौर पर, एक पत्थर को उस पर फेंक दें जिसने कहा था कि वीडियो कार्ड सशर्त संचालन के साथ खराब हैं: कृत्रिम रूप से लगभग 2 बार इस टुकड़े को धीमा करने में कामयाब रहे, लेकिन इसके लिए शर्तों की गहराई को 4 तक बढ़ाना आवश्यक था! if () if () if () if () if () ...

नैतिक:

• आधुनिक कार्डों पर, सामान्य रूप से, यह तर्क के साथ इतना बुरा नहीं है, लेकिन आपको नेस्टेड स्थितियों की एक बड़ी गहराई से बचना चाहिए।

जटिल डेटा संरचनाएं

इस विचार से प्रेरित कि एल्गोरिदम और डेटा संरचनाएं दृढ़ता से जुड़ी हुई हैं (कम से कम एन। Wirth), आपको यह देखना चाहिए कि कुछ जटिल डेटा संरचनाओं के साथ चीजें कैसे हैं।

यह वह जगह है जहां समस्या उत्पन्न होती है, डेटा को फ़ंक्शन में स्थानांतरित करते समय, हम केवल दो प्रकार की वस्तुओं का उपयोग कर सकते हैं - निरंतर अभिन्न प्रकार (अंक) और वीडियो मेमोरी ब्लॉकों के लिंक।

उदाहरण के लिए लिंक के आधार पर पेड़ों के निर्माण का विचार तुरंत शामिल किया गया है:

• हम कार्ड पर काम करने वाले फ़ंक्शन से मेमोरी आवंटित नहीं कर सकते हैं;
• किसी भी डेटा की किसी भी चयन और प्रतिलिपि बहुत धीमी है (अनुभाग 2 देखें)।

इस प्रकार, यह स्मृति के निरंतर ब्लॉक और इस ब्लॉक के तत्वों के संदर्भ में सरणी के रूप में जटिल डेटा संरचनाओं का प्रतिनिधित्व करने के लिए रहता है। तो आप आसानी से किसी भी डेटा ऐरे के ऊपर हैश टेबल, ट्री और इंडेक्स स्ट्रक्चर की कल्पना कर सकते हैं।

ऐसे ट्रिक्स के लिए पेबैक - डबल इंडेक्सिंग की आवश्यकता:

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
temp = data[index[j]+i];
}

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

आज हम बहुत कुछ समझ गए

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

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