ऐसा नहीं है कि बहुत पहले यह कंप्यूटिंग के लिए वीडियो कार्ड का उपयोग करने के लिए लोकप्रिय हो गया था। एक ठीक दिन, कुछ साल पहले, और मैंने नए, तब CUDA तकनीक को देखा। मेरे पास GTX8800 पर उस समय मेरे हाथों में एक अच्छा कार्ड था, और समानांतरकरण के कार्य भी थे।
जिसने भी GPU के साथ काम किया है वह क्वेरी पूलिंग, बैंक संघर्ष और उससे कैसे निपटना है, के बारे में जानता है, और यदि यह काम नहीं करता है, तो आप CUDA प्रोग्रामिंग [1] की मूल बातें पर कुछ उपयोगी लेख पा सकते हैं। GTX8800 कार्ड, एक अर्थ में, अच्छा था क्योंकि यह CUDA के पहले संस्करणों में से एक था और समर्थित था, इसलिए यह स्पष्ट रूप से उस पर दिखाई देता था जब बैंक संघर्ष होते हैं या वैश्विक मेमोरी के लिए अनुरोध संयुक्त नहीं होते हैं, क्योंकि इस मामले में समय बढ़ गया है बार। यह सब कार्ड के साथ काम करने और सामान्य कोड लिखने के सभी नियमों को बेहतर ढंग से समझने में मदद करता है।
नए मॉडल अधिक से अधिक कार्यक्षमता जोड़ रहे हैं, जो विकास को गति और गति प्रदान करता है। परमाणु संचालन, कैश, डायनेमिक कंसेंटर, आदि दिखाई दिए।
एक पोस्ट में मैं स्पोट-टेम्पोरल इमेज फ़िल्टरिंग और कम्प्यूटिंग क्षमता = 1.0 के लिए कार्यान्वयन के बारे में बात करूंगा, और नए फीचर्स के साथ परिणाम को गति दे सकता हूं।
सटीक पृष्ठभूमि के दमन की आवश्यकता होने पर उपग्रहों को देखने या अन्य फ़िल्टरिंग स्थितियों में अस्थायी फ़िल्टरिंग काम में आ सकती है।

सिद्धांत की बिट

जितना हो सके उतना कम किया, लेकिन सबसे महत्वपूर्ण बात छोड़ दी। तो, फ्रेम का एक क्रम इनपुट है। प्रत्येक फ्रेम को मैट्रिक्स, पृष्ठभूमि और उपयोगी सिग्नल के शोर के योग के रूप में दर्शाया जा सकता है।

कई पिछले फ़्रेमों का उपयोग करते हुए, आप नए आगमन वाले फ्रेम को सफेद कर सकते हैं और मौजूद होने पर सिग्नल का चयन कर सकते हैं। ऐसा करने के लिए, पिछले फ्रेम में कार्यान्वयन के अनुसार, इस फ्रेम पर पृष्ठभूमि की भविष्यवाणी करना, उसका अनुमान प्राप्त करना आवश्यक है। बता दें कि इतिहास में फ्रेम की संख्या N है, तो पृष्ठभूमि अनुमान को कुछ वज़न के साथ फ़्रेम के योग के रूप में दर्शाया जा सकता है।

गुणवत्ता मानदंड फ्रेम और अनुमानित पृष्ठभूमि के बीच का अंतर होगा, जिसे कम से कम किया जाना चाहिए।

समाधान का रूप है:

नमूना सहसंबंध मैट्रिक्स की गणना इस प्रकार की जाती है:

और दाईं ओर के वेक्टर का रूप है:

जहाँ k, l- 1 से N तक भिन्न होता है।
हालाँकि, यह विधि खाता फ़्रेम शिफ़्ट में नहीं ले जाती है। यदि आप गुणांक की कल्पना करते हैं

एक स्पष्ट बदलाव समारोह के रूप में

कुछ संदर्भ फ्रेम के सापेक्ष, फिर हम एक गुणवत्ता मानदंड तैयार कर सकते हैं और एक समाधान लिख सकते हैं। बहुत सारे सूत्र होने चाहिए, लेकिन परिणामस्वरूप

केवल पारियों पर निर्भर करता है।
कुछ संदर्भ फ्रेम के सापेक्ष परिवर्तन की गणना की जाती है।

निम्नलिखित सूत्रों के अनुसार:

संदर्भ फ्रेम के सापेक्ष फ्रेम शिफ्टिंग जानने के बाद, हम निम्न सूत्र के अनुसार गुणांक वेक्टर a (t) की गणना करते हैं:

जहां α डिग्री n के बहुपद का गुणांक वेक्टर है, is एक पर्याप्त छोटी मात्रा है, E पहचान मैट्रिक्स है।

GPU कार्यान्वयन

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

float **h_array_list, **d_array_list; h_array_list = ( float** )malloc( num_arrays * sizeof( float* )); cudaMalloc((void**)&d_array_list, num_arrays * sizeof( float * )); for ( int i = 0; i < num_arrays; i++ ) cudaMalloc((void**)&h_array_list[i], data_size); cudaMemcpy( d_array_list, h_array_list, num_arrays * sizeof(float*), cudaMemcpyHostToDevice );

इस तरह के एक संगठन के साथ, एक मामूली दोष है। यदि आप फ्रेम से निम्नानुसार पढ़ते हैं

 int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; int aBegin = nx * BLOCK_SIZE * by + BLOCK_SIZE * bx; float a = inFrame[ j ][ aBegin + nx * ty + tx]

फिर इस तरह के पढ़ने के साथ गणना क्षमता के पुराने संस्करणों में, वैश्विक मेमोरी के लिए अनुरोध संयुक्त नहीं थे।

यह अब एक कैश है और फ़्रेम को इंगित करने वाले इसमें जमा होते हैं और जल्दी से पढ़ते हैं। प्रश्नों के संयोजन के लिए नियमों को पूरा करने के लिए, आपको या तो पहले इन पॉइंटर्स को साझा मेमोरी (जो समय लगता है) में लोड करना होगा या, उदाहरण के लिए, यह करें:

 float **h_array_list, **d_array_list; h_array_list = ( float** )malloc( 16 * num_arrays * sizeof( float* ) ); cudaMalloc((void**)&d_array_list, 16 * num_arrays * sizeof(float *)); for ( int i = 0; i < num_arrays; i++ ) { cudaMalloc((void**)&h_array_list[i*16], data_size); for( int j = 1; j < 16; j++ ) h_array_list[ i*16 + j ] = h_array_list[ i*16 ]; } cudaMemcpy( d_array_list, h_array_list, 16*num_arrays * sizeof(float*), cudaMemcpyHostToDevice );

इस मामले में, पढ़ने का अनुरोध संयुक्त होगा और हम गति में एक उल्लेखनीय वृद्धि को नोटिस करेंगे।
GPU के लिए फ्रेम को सफेद करने का कार्य कुछ इस तरह दिखता है:

 __global__ void kernelWhite( float * out, float ** inFrame, float *at, int nx, int mem ) { int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; int aBegin = nx * BLOCK_SIZE * by + BLOCK_SIZE * bx; __shared__ float frame1 [BLOCK_SIZE][BLOCK_SIZE]; __shared__ float a [32]; int i = ty * BLOCK_SIZE + tx; //    if( i < mem ) a[ i ] = at[ i ]; __syncthreads (); frame1 [ty][tx] = 0; //         for( int j = 0; j < mem; j++ ) { frame1 [ty][tx] += a[ j ]*inFrame[ j*16 + tx ][ aBegin + nx * ty + tx]; } //   out [ aBegin + nx * ty + tx ] = inFrame[ mem*16 + tx ][ aBegin + nx * ty + tx] - frame1 [ty][tx]; }

सामान्य तौर पर, यहां सब कुछ सरल है, हम फ्रेम को सफेद करने के परिणाम के लिए एक साझा मेमोरी बनाते हैं, हमारे गुणांक वेक्टर को लोड करते हैं, प्रवाह को सिंक्रनाइज़ करते हैं, और फिर एक चक्र में इन गुणांक के साथ हमारे फ्रेम को साइकिल करते हैं। यहां आप प्रवाह के एक वर्ग ब्लॉक को भी नहीं चुन सकते हैं, लेकिन एक रैखिक एक, जो सामान्य रूप से और भी अधिक सही होगा, क्योंकि यह अनुक्रमण में गणना की संख्या को कम करेगा।
शिफ्ट खोज फ़ंक्शन के लिए, एक मानक अनुकूलित एल्गोरिथ्म का उपयोग सरणी [2] के योग को खोजने के लिए किया गया था। इस फ़ंक्शन में, आपको पहले गुणांक रिज, जीआईज, हिज को ढूंढना होगा और फिर उपरोक्त सूत्रों का उपयोग करके योगों की गणना करना होगा।

__global__ शून्य कर्नेलशफ़्ट (फ़्लोट * इनफ्रैम, फ़्लोट * इनफ़्रैम, फ़्लोट * आउटडेटा, इंट एनएक्स)

 __global__ void kernelShift ( float *inFrameR, float *inFrame, float *outData, int nx ) { int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; int aBegin = nx * BLOCK_SIZE * by + BLOCK_SIZE * bx; //        __shared__ float rij [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float gij [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float hij [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float tmp [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float p [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float q [BLOCK_SIZE*BLOCK_SIZE]; __shared__ float frame1 [BLOCK_SIZE][BLOCK_SIZE]; //    shared memeory frame1 [ty][tx] = inFrameR [ aBegin + nx * ty + tx]; //  __syncthreads (); int i = ty * BLOCK_SIZE + tx; //   gij, hij, rij   1414. if( tx > 0 && > 0 && tx < BLOCK_SIZEm1 && ty < BLOCK_SIZEm1 ) { gij[ i ] = ( frame1[ty + 1][tx] - frame1[ty - 1][tx] )/2.0; hij[ i ] = ( frame1[ty][tx + 1] - frame1[ty][tx - 1] )/2.0; rij[ i ] = ( frame1[ty][tx] + frame1[ty][tx + 1] + frame1[ty][tx - 1] + frame1[ty + 1][tx] + frame1[ty - 1][tx] )/5.0; } else { gij[ i ] = 0; hij[ i ] = 0; rij[ i ] = 0; } //  __syncthreads (); //    shared memeory frame1 [ty][tx] = inFrame [ aBegin + nx * ty + tx] - rij[i]; //  p, q p [i] = gij [i]*frame1 [ty][tx]; q [i] = hij [i]*frame1 [ty][tx]; //  u tmp [i] = hij [i]*hij [i]; //  v rij[i] = hij [i]*gij [i]; //  w hij [i] = gij [i]*gij [i]; __syncthreads (); for ( int s = BLOCK_SIZEhalfqrt; s > 32; s >>= 1 ) { if ( i < s ) { p [i] += p [i + s]; q [i] += q [i + s]; tmp [i] += tmp [i + s]; rij [i] += rij [i + s]; hij [i] += hij [i + s]; } __syncthreads (); } if ( i < 32 ) { p [i] += p [i + 32]; p [i] += p [i + 16]; p [i] += p [i + 8]; p [i] += p [i + 4]; p [i] += p [i + 2]; p [i] += p [i + 1]; q [i] += q [i + 32]; q [i] += q [i + 16]; q [i] += q [i + 8]; q [i] += q [i + 4]; q [i] += q [i + 2]; q [i] += q [i + 1]; tmp [i] += tmp [i + 32]; tmp [i] += tmp [i + 16]; tmp [i] += tmp [i + 8]; tmp [i] += tmp [i + 4]; tmp [i] += tmp [i + 2]; tmp [i] += tmp [i + 1]; rij [i] += rij [i + 32]; rij [i] += rij [i + 16]; rij [i] += rij [i + 8]; rij [i] += rij [i + 4]; rij [i] += rij [i + 2]; rij [i] += rij [i + 1]; hij [i] += hij [i + 32]; hij [i] += hij [i + 16]; hij [i] += hij [i + 8]; hij [i] += hij [i + 4]; hij [i] += hij [i + 2]; hij [i] += hij [i + 1]; } if ( i == 0 ) { outData [ by*gridDim.x + bx ] = p[0]; outData [ by*gridDim.x + bx + gridDim.x*gridDim.y] = q[0]; outData [ by*gridDim.x + bx + 2*gridDim.x*gridDim.y] = tmp [0]; outData [ by*gridDim.x + bx + 3*gridDim.x*gridDim.y] = rij [0]; outData [ by*gridDim.x + bx + 4*gridDim.x*gridDim.y] = hij [0]; } };

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

 //    . int numbCoeff = regInfo->nx / threads.x*regInfo->nx / threads.y; cudaMemcpy( coeffhost, coeffdev, numbCoeff*5* sizeof ( float ), cudaMemcpyDeviceToHost ); float q = 0, p = 0; float u = 0, v = 0, w = 0; for( int i = 0; i < numbCoeff; i++ ) { p += coeffhost[ i ]; q += coeffhost[ i + numbCoeff]; u += coeffhost[ i + 2*numbCoeff]; v += coeffhost[ i + 3*numbCoeff]; w += coeffhost[ i + 4*numbCoeff]; } dx = ( q * w - p * v )/( u * w - v * v ); dy = ( p * u - q * v )/( u * w - v * v );

यहां मुख्य नुकसान क्या हैं।
1. अंतिम परिणाम हमें प्रोसेसर पर मिलता है।
2. संदर्भ फ्रेम रिज, जीआईजे, हिज के लिए गुणांक प्रत्येक फ्रेम के लिए बार-बार गणना की जाती है।

आइए प्रोफाइलर में परिणाम देखें और मूल्यांकन करें कि क्या यह कार्यान्वयन अच्छी तरह से काम करता है। प्रोफाइलर हमें बताता है कि पढ़ने के अनुरोध संयुक्त हैं, और निर्देश अधिकांश समय लेते हैं।
ऑफसेट खोज कर्नेल:
ऑफसेट खोज कोर।

सफेद फ्रेम के लिए कर्नेल:
वाइटनिंग फ्रेम के लिए कर्नेल।

उत्थान और अनुकूलन

चूंकि पुराना GTX8800 एक अलग कंप्यूटर पर था और पहले से ही कहीं चला गया था, इसलिए मैं GTX टाइटन के लिए परिणाम देता हूं।
512x512 आकार के फ्रेम के लिए, निम्नलिखित कार्य समय परिणाम प्राप्त किए गए थे। कतरनी गुणांक की गणना के लिए कोर 60 μs लेता है। अगर हम इसे गुणांक और अंतिम योग की प्रतिलिपि में जोड़ते हैं, तो हमें 130 μ मिलते हैं। 20 फ्रेम की मेमोरी के साथ, ऑपरेटिंग समय पहले से ही क्रमशः 2.6 एमएस है। वाइटनिंग फ्रेम का कार्य 156 μs लेता है। प्रोसेसर पर गुणांक की गणना करने के लिए एक फ़ंक्शन भी है, यह एक और 50 μ है। कुल में, 2.806 एमएस प्राप्त किया जाता है।

तदनुसार, सीपीयू पर, पारियों की गणना में 40 एमएस लगता है, फ्रेम 30 एमएस की सफेदी और अन्य 50 μs के गुणांक की गणना होती है। कुल में, यह 70 एमएस निकला, जो कि जीपीयू की तुलना में 25 गुना धीमा है (वास्तव में, ऐसा नहीं है, लेकिन बाद में उस पर अधिक)।

कार्यान्वयन एसएस = 1.0 के साथ काफी सरल और संगत है। आइए अब देखें कि क्या तय किया जा सकता है। पारियों की खोज में अंतिम क्रिया को परमाणु परिचालनों के माध्यम से बदला और कार्यान्वित किया जा सकता है, जो कई गुना तेजी से काम करेगा।

संशोधित कोड अब इस तरह दिखता है:

 //       if ( i == 0 ) { atomicAdd( &outData[0], p[0] ); atomicAdd( &outData[1], q[0] ); atomicAdd( &outData[2], tmp[0] ); atomicAdd( &outData[3], rij[0] ); atomicAdd( &outData[4], hij[0] ); } //  CPU   float p = coeffhost[ 0 ]; float q = coeffhost[ 1 ]; float u = coeffhost[ 2 ]; float v = coeffhost[ 3 ]; float w = coeffhost[ 4 ]; dx = ( q * w - p * v )/( u * w - v * v ); dy = ( p * u - q * v )/( u * w - v * v );

इस मामले में, कर्नेल लगभग एक ही समय, 66 μs लेता है।

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

GPU पर हमारे कार्य केवल समय की एक छोटी अवधि में काम करते हैं, बाकी की देरी, मेमोरी और अन्य मध्यवर्ती संचालन की नकल करते हैं। कुल मिलाकर, 11.5 एमएस सभी फ्रेम के लिए बदलाव की गणना करने के लिए चलता है। परिणामस्वरूप, सीपीयू के सापेक्ष त्वरण केवल 6 गुना था। हम सभी मध्यवर्ती गणनाओं को हटाते हैं और गुठली को एक के बाद एक लूप में चलाते हैं:

 for( int j = 0; j < regressInfo->memory - 1; j++ ) kernelShift<<< SKblocks, SKthreads >>> ( … );

अनुकूलन के बाद, प्रोफाइलर एक पूरी तरह से अलग तस्वीर दिखाता है।

बहुत अच्छा हुआ, ठीक है? .. अब पाली की गणना के साथ हिस्सा 11.5 एमएस के बजाय 1.40 एमएस लेता है। 4 μ के बराबर कोर के बीच एक छोटा अंतराल बना रहा, जो एक प्रोफाइल बनाता है। और अब चलो विभिन्न धागे में गुठली को चलाने की कोशिश करते हैं।

 // stream for( int it = 0; it < 4; it++ ) cudaStreamCreate( &streamTX[it] ); …… for( int j = 0; j < regressInfo->memory - 1; j++ ) kernelShift<<< SKblocks, SKthreads, 0, streamTX[j%4] >>> (…);

अब प्रोफाइलर ऐसा परिणाम दिखाता है और यहाँ अंतराल 1.22 एमएस है।

और यह 0.66 एमएस से 0.066 * 20 = 1.320 से भी कम है।
प्रोसेसर लोड 84% है।

नतीजतन, कुल समय 1.22ms + 50 μs + 156 μs = 1.426 ms है, जो CPU की तुलना में 50 गुना तेज है।
बेशक, बदलावों की गणना के लिए एक कोर लिखना अधिक सही होगा, लेकिन यह एक अलग अनुकूलन मार्ग है, और हमने अधिक सामान्य तरीका माना है, जिसे कम करके नहीं आंका जा सकता है। आप GPU पर एक अच्छा कार्यान्वयन लिख सकते हैं, लेकिन GPU के साथ सहभागिता को अनुकूलित करने के बारे में भूल जाते हैं। कर्नेल को जितना कम समय लगेगा, कॉल के बीच का अंतराल उतना अधिक होगा जिसमें सीपीयू पर कोई क्रिया होती है।

फ़िल्टर परिणाम

उदाहरण के लिए, तारों वाले आकाश की तस्वीरें ली गईं, एक मंद मंद वस्तु, छोटे कंपन और आंतरिक शोर को जोड़ा गया।
प्रसंस्करण से पहले:

प्रसंस्करण के बाद:

प्रसंस्करण से पहले:

प्रसंस्करण के बाद:

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

PS मुझे एक वाइड-एंगल कैमरा से 4GB 10Mpix फ्रेम मिला। मैं इसे प्रोसेस करने की कोशिश करूंगा।

1. CUDA में मेमोरी के साथ काम करें।
2. CUDA पर राशि की गणना के लिए एल्गोरिथ्म।

GPU पर छवियों के स्थानिक-लौकिक प्रसंस्करण

सिद्धांत की बिट

GPU कार्यान्वयन

उत्थान और अनुकूलन

फ़िल्टर परिणाम

More articles: