छवियों का सुपरफास्ट मार्कअप

लेख में मैं आपको बताऊंगा कि कैसे बहुत जल्दी से एक द्विआधारी रेखापुंज पर जुड़ी वस्तुओं को सूचीबद्ध करें, जो मैंने पिछले लेख में कहा था। ऐसा प्रतीत होगा कि ऐसी गति कहाँ है; अब हम दसियों मिलीसेकंड के लिए 4096 तक 4096 चित्रों के साथ "टूटेंगे"। और यद्यपि यह कार्य अपने आप में दिलचस्प है, लेकिन इसका समाधान काफी व्यापक प्रयोज्यता के साथ सरल और मूल विधि पर आधारित है, जिसका मुख्य शोध "हम इसे सरल बनाएंगे और देखेंगे कि इसमें क्या आता है"। इस मामले में, CUDA का उपयोग मुख्य कैलकुलेटर के रूप में किया जाएगा, लेकिन बिना किसी बारीकियों के, क्योंकि हम इसे "बहुत सरल" बनाना चाहते हैं।

कार्य

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

उदाहरण के लिए, मूल चित्र:


एल्गोरिथ्म बहुत बड़ी छवियों (16384x16384 तक) के साथ काम करेगा, लेकिन स्पष्ट कारणों के लिए, लेख के पाठ में कम छवियों का उपयोग किया जाएगा।

अन्वेषण करें या नहीं?

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

सबसे खराब "सबसे अच्छा मामला"

जाहिर है, समस्या को हल करने के लिए, आपको छवि के प्रत्येक पिक्सेल पर विचार करने की आवश्यकता है। देखते हैं कि मॉडल कोड 8192x8192 की छवि पर कब तक चलेगा:

for ( int x = 0; x < dims[0]; x++)
{
for ( int y = 0; y < dims[1]; y++)
{
image[x + dims[0]*y] = 0;
}
}


हां, चक्र को एक में बदलकर इसे "अनुकूलित" किया जा सकता है, लेकिन किसी भी उचित विधि में छवि के व्यक्तिगत पिक्सेल तक पहुंच की संभावना शामिल है, ताकि अनुक्रमण जानबूझकर छोड़ दिया जाए। इसे कुछ "औसत" प्रोसेसर (मेरे मामले में, कोर 2 डुओ E8400) पर चलाएं और लगभग 680-780 एमएस का रनटाइम प्राप्त करें।

यह कहना मुश्किल है कि यह बहुत कम है या थोड़ा है, लेकिन ऐसा लगता है कि समस्या को तेजी से हल करना संभव नहीं होगा।

या इसलिए यह सिर्फ लगता है; क्योंकि "सभी साधन अच्छे हैं" और ऐसी प्रौद्योगिकियाँ हैं जो समस्या के समाधान को बहुत सरल करती हैं और बेहतर परिणाम देती हैं; और इस मामले में, CUDA तकनीक, जिसमें गहनता से मेमोरी का उपयोग करने वाले कार्यों के लिए बहुत अच्छा प्रदर्शन है, हमारी बहुत मदद कर सकता है।

मेरे औसत GTX460 ग्राफिक्स कार्ड पर, हमें मुख्य मेमोरी में लगभग 40 जीबी / एस बनाम 2 जीबी / एस की वीडियो मेमोरी तक पहुंच मिलती है। यह प्रसंस्करण गति बड़े पैमाने पर समानता और काफी तेज मेमोरी के उपयोग के माध्यम से हासिल की जाती है।

फूट डालो और जीतो

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

मूल छवि को विभाजित करने का एक उचित तरीका इसे छोटे "वर्गों" में विभाजित करना प्रतीत होता है, जिनमें से प्रत्येक को एक अलग स्ट्रीम में संसाधित किया जाएगा; यह CUDA गणना मॉडल में अच्छी तरह से फिट बैठता है। और यदि आप अत्यधिक विशिष्टता को छोड़ देते हैं (जो एक परिवर्तन द्वारा किया जाता है):

UID = threadIdx.x + blockDim.x * blockIdx.x

तब हमें पता चलता है कि मूल कार्य को प्रवाह (यूआईडी) के सभी विशिष्ट पहचानकर्ताओं के लिए उपकथाओं के परिणामों के संघ के रूप में दर्शाया गया है। हां, हम अभी भी नहीं जानते हैं कि हम वास्तव में क्या करेंगे, लेकिन हम पहले से ही जानते हैं कि कैसे कार्य को तोड़ना है।

इसलिए (रंगों का दंगा लेखक की शैली है, और ढाल से पता चलता है कि सभी यूआईडी अलग हैं) हम मूल छवि को कवर कर सकते हैं:


प्रत्येक स्ट्रीम पिक्सल के अनुरूप ब्लॉक में प्रक्रिया करेगी, उदाहरण के लिए, आकार में 32 x 32। फिर, कुछ तुल्यकालन और परिणामों के संयोजन की आवश्यकता हो सकती है।

जितना सरल है

गणित में, एक अद्भुत सिद्धांत अक्सर उपयोग किया जाता है: "यदि कुछ असंभव है, लेकिन आप वास्तव में चाहते हैं, तो इसका मतलब है कि आप कर सकते हैं", वास्तव में यह "बाद की जटिलताओं" को "पीछे" धकेल देता है, लेकिन इस मामले में यह काम में आएगा।

हम सरलतम फ़िल्ट एल्गोरिथ्म लेते हैं और प्रत्येक ब्लॉक के भीतर, सीमा के प्रभावों के बारे में सोचने के बिना, इसे संबंधित बिंदुओं पर लागू करते हैं (प्रत्येक से अभी तक भरा नहीं है)।

वास्तव में, हम उन छवि बिंदुओं को भरते हैं जो अपने विशिष्ट पहचानकर्ता के साथ ब्लॉक में हैं, परिणाम अनुमानित है:



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



यह समझना आसान है कि हम तस्वीर में "ग्रिड" के साथ जीते हैं जो संसाधित वर्गों की सीमाओं का प्रतिनिधित्व करते हैं:



अब, प्रत्येक बॉक्स के भीतर, जुड़े हुए क्षेत्र एक रंग से भरे हुए हैं, डिस्कनेक्ट किए गए अलग हैं। यह प्रत्येक बॉक्स के लिए O (n ^ 3) के लिए काम करेगा (हां, मैंने फिल एल्गोरिथ्म का सबसे खराब कार्यान्वयन चुना), जिनमें से (imageSize.x / n) * (imageSize.y / n) हैं, जिनमें से अधिकांश को एक साथ पर्याप्त रूप से संसाधित किया जाएगा। धागे की संख्या (आमतौर पर लगभग 256)।

आस्थगित कठिनाइयों

यह परिणामों को एक में एकत्रित करने के लिए बना हुआ है। इस तथ्य के बावजूद कि जागरूकता के स्तर पर, उपकेंद्रों के लिए शास्त्रीय पद्धति को लागू करने की तुलना में अधिक कठिन है, एल्गोरिदमिक रूप से यह बहुत तेज है।

वास्तव में, हम यह समझना चाहते हैं कि विभिन्न रंग एक ही जुड़े क्षेत्र के हैं।

ऐसा करने के लिए, बस ग्रिड के पिक्सेल के माध्यम से जाएं (चित्र में बिल्कुल सफेद रंग में खींचा गया है) और जांचें कि क्या अलग-अलग रंगों के पिक्सल तत्काल पड़ोसी हैं। यदि वे हैं, तो दोनों रंगों का अर्थ अनिवार्य रूप से एक वस्तु है जिसे हम आसानी से "पुनरावृत्ति" कर सकते हैं, या इस तथ्य को ध्यान में रख सकते हैं (आमतौर पर वस्तुओं के मापदंडों को प्राप्त करने के लिए पर्याप्त)।

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

इन बिंदुओं के माध्यम से जाने के बाद, आप एक रंग मिलान मानचित्र (c ₁ == c ₂ == c ₃ ...) खींच सकते हैं और परिणाम को ठीक कर सकते हैं:

कार्यान्वयन और गति का अनुमान

हैरानी की बात है, इस पद्धति का "अड़चन" मुख्य मेमोरी से वीडियो मेमोरी में डेटा का हस्तांतरण है, यहां हमारे पास कोई विकल्प नहीं है:

START_ACTION( "Allocating GPU memory: image... " );
unsigned int * data_cuda;
cutilSafeCall( cudaMalloc(( void **)&data_cuda, data_cuda_size) );
END_ACTION(data_cuda_size);

START_ACTION( "Copying input data to GPU memory: image... " );
cutilSafeCall( cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice) );
END_ACTION(data_cuda_size);


परिणामी 2 जीबी / एस आश्चर्य की बात नहीं है - यह रैम तक पहुंच की गति है। जिन छवियों के साथ हम काम करते हैं - ये मिलीसेकंड हैं, डरावने नहीं।

कर्नेल को कॉल करना (वीडियो कार्ड पर किया जाने वाला कार्य):

START_ACTION( "Calling CUDA kernel... " );
cutilSafeCall( cudaThreadSynchronize() );
processKernelCUDA<<<GRID, THREADS>>>(data_cuda, dims[0], dims[1]);
cutilSafeCall( cudaThreadSynchronize() );
END_ACTION(data_cuda_size);


निष्कर्ष खुद के लिए बोलता है:

Calling CUDA kernel... 26.377720 msecs.
Throughput = 2.36942 GBytes/s

वास्तव में, वीडियो मेमोरी में आवश्यक डेटा (छवि) के कम हस्तांतरण के समय के दौरान, सभी उप-प्रकारों की पूर्ति का एहसास करना संभव था। यह इस तथ्य के बावजूद है कि मैंने सबसे खराब भरण एल्गोरिथ्म चुना, और इसके कार्यान्वयन को बिल्कुल इष्टतम नहीं बनाया। क्या यह कुछ का अनुकूलन करने के लिए समझ में आता है? आप तय करें :)

कर्नेल कोड के बिना, कथन संभवतः अधूरा रहेगा, इसलिए कृपया प्यार और अनुग्रह करें:

__global__ void processKernelCUDA(unsigned int * image, int dim_x, int dim_y)
{
// Full task size: image dimensions
int TASK_SIZE = dim_x * dim_y;

// Max UID for full task (really count of 'quads')
int MAX_UID = TASK_SIZE / QUAD_SIZE / QUAD_SIZE;

// calculate UID by CUDA implicit parameters; WORKERS = GRID*THREADS
for ( int UID = threadIdx.x + blockDim.x * blockIdx.x; UID < MAX_UID; UID += WORKERS)
{
// count of quads for x-dimension
int quads_x = dim_x / QUAD_SIZE;

// current quad index form UID
int curr_quad_x = UID % quads_x;
int curr_quad_y = UID / quads_x;

// offset in global task
int offset = curr_quad_x * QUAD_SIZE + dim_x * curr_quad_y * QUAD_SIZE;

// id of 'color' to mark block, have to be >= 2
int COLOR_ID = 2 + UID;

do
{
// means that fill is not started yet
bool first_shot = true ;

// process quad starting with specified 'offset'
bool done_something;

do // fill step
{
done_something = false ;
for ( int local_y = 0; local_y < QUAD_SIZE; local_y++) // step by y
{
// calcluate point index in whole image
int start_x = offset + local_y * dim_x;
for ( int pos = start_x; pos < start_x + QUAD_SIZE; pos++) // step by x
{
if (image[pos] == 1) // is not marked with ID, is not blank
{
if ( first_shot /* start fill */ // else check neighbors:
|| image[pos-1] == COLOR_ID // x--
|| image[pos+1] == COLOR_ID // x++
|| local_y > 0 && image[pos-dim_x] == COLOR_ID // y--
|| local_y < QUAD_SIZE && image[pos+dim_x] == COLOR_ID ) // y++
{
image[pos] = COLOR_ID; // mark with ID
first_shot = false ; // fill already started
done_something = true ; // fill step does something
}
}
}
}
} while (done_something); // break if fill have no results

if (first_shot)
{
// no more work in this block, exit
break ;
}
else
{
// next iteration: restart block with different color ID
COLOR_ID += MAX_UID;
}
} while ( true );
}
}



इनपुट तस्वीर के लिए एक संकेतक है और एक्स और वाई के अंदर तस्वीर का रिज़ॉल्यूशन है। सबसैक आइडेंटिफ़ायर की गणना की जाती है, सबस्कॉक के अनुरूप वर्ग की ऑफसेट को चुना जाता है, और यह एक गैर-पुनरावर्ती एल्गोरिथ्म से भरा होना शुरू होता है।

सफलता का राज

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

पुनश्च एक बड़ा अनुरोध, चित्रों को कॉपी न करें और मेरी अनुमति के बिना फिर से प्रकाशित न करें (मैं विधि का लेखक हूं)।